llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Akimasa Watanuki (Men-cotton) <details> <summary>Changes</summary> 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 --- Full diff: https://github.com/llvm/llvm-project/pull/219688.diff 4 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+13) - (modified) clang/lib/CIR/CodeGen/CIRGenModule.h (+1) - (added) clang/test/CIR/CodeGenOpenCL/version.cl (+15) - (modified) clang/test/CodeGenCUDASPIRV/kernel-cc.cu (+4) ``````````diff 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> `````````` </details> https://github.com/llvm/llvm-project/pull/219688 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
