Author: Akimasa Watanuki Date: 2026-09-12T13:26:30+09:00 New Revision: ba3010d7e4520b6aee926f878c5c1771f1c2a6f3
URL: https://github.com/llvm/llvm-project/commit/ba3010d7e4520b6aee926f878c5c1771f1c2a6f3 DIFF: https://github.com/llvm/llvm-project/commit/ba3010d7e4520b6aee926f878c5c1771f1c2a6f3.diff LOG: [CIR][OpenCL] Emit OpenCL language version metadata in CIR (#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 Added: clang/test/CIR/CodeGenOpenCL/version.cl clang/test/SemaHIP/atomic-init.hip Modified: clang/lib/CIR/CodeGen/CIRGenModule.cpp clang/lib/CIR/CodeGen/CIRGenModule.h clang/lib/CodeGen/CodeGenModule.cpp clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip clang/test/CodeGenCUDASPIRV/kernel-cc.cu clang/test/CodeGenHIP/hipspv-kernel.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 0dcd35269dbdc..6da830289568e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -136,6 +136,15 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, theModule->setAttr( cir::CIRDialect::getSourceLanguageAttrName(), cir::SourceLanguageAttr::get(&mlirContext, *sourceLanguage)); + if (langOpts.OpenCL || (langOpts.CUDAIsDevice && getTriple().isSPIRV())) { + // CUDA and HIP use OpenCL 2.0 metadata when targeting SPIR-V. + unsigned version = + langOpts.OpenCL ? langOpts.getOpenCLCompatibleVersion() : 200; + setOpenCLVersionAttr(cir::CIRDialect::getOpenCLVersionAttrName(), version); + 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 @@ -199,6 +208,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 5646db9503dc6..51b9c420c94be 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/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index f3d1524ce94e6..40feb4b30f98c 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2014,7 +2014,9 @@ void CodeGenModule::EmitOpenCLMetadata() { // SPIR v2.0 s2.13 - The OpenCL version used by the module is stored in the // opencl.ocl.version named metadata node. // C++ for OpenCL has a distinct mapping for versions compatible with OpenCL. - auto CLVersion = LangOpts.getOpenCLCompatibleVersion(); + // CUDA and HIP use OpenCL 2.0 metadata when targeting SPIR-V. + unsigned CLVersion = + LangOpts.OpenCL ? LangOpts.getOpenCLCompatibleVersion() : 200; auto EmitVersion = [this](StringRef MDName, int Version) { llvm::Metadata *OCLVerElts[] = { diff --git a/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip index 039ab35f1c906..45483ab5790cb 100644 --- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip +++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip @@ -10,6 +10,8 @@ // Test that HIP kernels on AMDGCN-flavored SPIR-V get the spir_kernel // calling convention. +// CIR: cir.cl.version = #cir.cl.version<2, 0> + #define __global__ __attribute__((global)) #define __device__ __attribute__((device)) 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> diff --git a/clang/test/CodeGenHIP/hipspv-kernel.cpp b/clang/test/CodeGenHIP/hipspv-kernel.cpp index b9a6df12c7968..1ef8a5b19ccfb 100644 --- a/clang/test/CodeGenHIP/hipspv-kernel.cpp +++ b/clang/test/CodeGenHIP/hipspv-kernel.cpp @@ -1,5 +1,9 @@ // RUN: %clang_cc1 -triple spirv64 -x hip -emit-llvm -fcuda-is-device \ // RUN: -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -emit-llvm -fcuda-is-device \ +// RUN: -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device \ +// RUN: -o - %s | FileCheck %s --check-prefix=AMDGCN #define __global__ __attribute__((global)) @@ -7,3 +11,7 @@ __global__ void foo(float *a, float b) { *a = b; } + +// CHECK: !opencl.ocl.version = !{[[OCL:![0-9]+]]} +// CHECK: [[OCL]] = !{i32 2, i32 0} +// AMDGCN-NOT: !opencl.ocl.version diff --git a/clang/test/SemaHIP/atomic-init.hip b/clang/test/SemaHIP/atomic-init.hip new file mode 100644 index 0000000000000..6bf9026417484 --- /dev/null +++ b/clang/test/SemaHIP/atomic-init.hip @@ -0,0 +1,8 @@ +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fcuda-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -fsyntax-only -verify %s +// expected-no-diagnostics + +// SPIR-V's OpenCL metadata must not impose OpenCL initialization restrictions. +__attribute__((device)) void atomic_init() { + _Atomic(int) x = 0; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
