https://github.com/Men-cotton created https://github.com/llvm/llvm-project/pull/219688
Emit OpenCL and C++ for OpenCL language version attributes from CIRGen. Preserve the compatible OpenCL version and the C++ for OpenCL version separately so later lowering does not infer one from the other. Assisted-by: Codex / GPT-5.6 Sol >From 8a4e9381a77667a5ab2dcd824cf2f1b3ae94baff Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Mon, 8 Jun 2026 20:56:59 +0900 Subject: [PATCH] [CIR][OpenCL] Emit OpenCL language version metadata in CIR Emit OpenCL and C++ for OpenCL language version attributes from CIRGen. Preserve the compatible OpenCL version and the C++ for OpenCL version separately so later lowering does not infer one from the other. Assisted-by: Codex / GPT-5.6 Sol --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 13 +++++++++++++ clang/lib/CIR/CodeGen/CIRGenModule.h | 1 + clang/test/CIR/CodeGenOpenCL/version.cl | 15 +++++++++++++++ clang/test/CodeGenCUDASPIRV/kernel-cc.cu | 4 ++++ 4 files changed, 33 insertions(+) create mode 100644 clang/test/CIR/CodeGenOpenCL/version.cl diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index d61743de3a8e4..bd505621c143a 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -135,6 +135,13 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, theModule->setAttr( cir::CIRDialect::getSourceLanguageAttrName(), cir::SourceLanguageAttr::get(&mlirContext, *sourceLanguage)); + if (langOpts.OpenCL || (langOpts.CUDAIsDevice && getTriple().isSPIRV())) { + setOpenCLVersionAttr(cir::CIRDialect::getOpenCLVersionAttrName(), + langOpts.getOpenCLCompatibleVersion()); + if (langOpts.OpenCLCPlusPlus) + setOpenCLVersionAttr(cir::CIRDialect::getOpenCLCXXVersionAttrName(), + langOpts.OpenCLCPlusPlusVersion); + } theModule->setAttr(cir::CIRDialect::getTripleAttrName(), builder.getStringAttr(getTriple().str())); // TODO(CIR): These attributes should eventually be replaced by @@ -198,6 +205,12 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, CIRGenModule::~CIRGenModule() = default; +void CIRGenModule::setOpenCLVersionAttr(StringRef attrName, unsigned version) { + theModule->setAttr( + attrName, cir::OpenCLVersionAttr::get(&getMLIRContext(), version / 100, + (version % 100) / 10)); +} + void CIRGenModule::createCUDARuntime() { cudaRuntime.reset(createNVCUDARuntime(*this)); } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index 012478cc2ad23..e0057cfe8d4de 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -140,6 +140,7 @@ class CIRGenModule : public CIRGenTypeCache { void createCUDARuntime(); void createOpenMPRuntime(); + void setOpenCLVersionAttr(llvm::StringRef attrName, unsigned version); /// A helper for constructAttributeList that handles return attributes. void constructFunctionReturnAttributes(const CIRGenFunctionInfo &info, diff --git a/clang/test/CIR/CodeGenOpenCL/version.cl b/clang/test/CIR/CodeGenOpenCL/version.cl new file mode 100644 index 0000000000000..636f52e676604 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/version.cl @@ -0,0 +1,15 @@ +// RUN: %clang_cc1 -cl-std=CL1.2 -fclangir -emit-cir -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CL12-CIR %s +// RUN: %clang_cc1 -cl-std=CL3.0 -fclangir -emit-cir -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CL30-CIR %s +// RUN: %clang_cc1 -x clcpp -cl-std=CLC++ -fclangir -emit-cir -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CLCXX10-CIR %s +// RUN: %clang_cc1 -x clcpp -cl-std=CLC++2021 -fclangir -emit-cir -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CLCXX2021-CIR %s + +// CL12-CIR: cir.cl.version = #cir.cl.version<1, 2> +// CL30-CIR: cir.cl.version = #cir.cl.version<3, 0> +// CLCXX10-CIR-DAG: cir.cl.cxx.version = #cir.cl.version<1, 0> +// CLCXX10-CIR-DAG: cir.cl.version = #cir.cl.version<2, 0> +// CLCXX2021-CIR-DAG: cir.cl.cxx.version = #cir.cl.version<2021, 0> +// CLCXX2021-CIR-DAG: cir.cl.version = #cir.cl.version<3, 0> + +__kernel void version_marker(__global int *out) { + out[0] = 1; +} diff --git a/clang/test/CodeGenCUDASPIRV/kernel-cc.cu b/clang/test/CodeGenCUDASPIRV/kernel-cc.cu index 9e575d232b34d..a525b4077ef87 100644 --- a/clang/test/CodeGenCUDASPIRV/kernel-cc.cu +++ b/clang/test/CodeGenCUDASPIRV/kernel-cc.cu @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -fcuda-is-device -triple spirv32 -o - -emit-llvm -x cuda %s | FileCheck %s // RUN: %clang_cc1 -fcuda-is-device -triple spirv64 -o - -emit-llvm -x cuda %s | FileCheck %s +// RUN: %if cir-enabled %{ %clang_cc1 -fcuda-is-device -triple spirv32 -o - -emit-cir -fclangir -x cuda %s | FileCheck %s --check-prefix=CIR %} +// RUN: %if cir-enabled %{ %clang_cc1 -fcuda-is-device -triple spirv64 -o - -emit-cir -fclangir -x cuda %s | FileCheck %s --check-prefix=CIR %} // Verifies that building CUDA targeting SPIR-V {32,64} generates LLVM IR with // spir_kernel attributes for kernel functions. @@ -10,3 +12,5 @@ __attribute__((global)) void kernel() { return; } // CHECK: !opencl.ocl.version = !{[[OCL:![0-9]+]]} // CHECK: [[OCL]] = !{i32 2, i32 0} + +// CIR: cir.cl.version = #cir.cl.version<2, 0> _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
