https://github.com/Men-cotton updated https://github.com/llvm/llvm-project/pull/219688
>From b37e6db449417bee7c9342939f33008d1e0901d0 Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Mon, 8 Jun 2026 20:56:59 +0900 Subject: [PATCH 1/2] [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 0dcd35269dbdc1..322b0327fa3abd 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -136,6 +136,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 @@ -199,6 +206,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 5646db9503dc63..51b9c420c94bed 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 00000000000000..636f52e676604d --- /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 9e575d232b34de..a525b4077ef87e 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> >From 871a42d957443c6cc242408d4ad8d976b9f6cb6e Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Sat, 5 Sep 2026 22:31:47 +0900 Subject: [PATCH 2/2] fix: Supply the HIP SPIR-V version only for metadata emission Fix the assertion exposed by PR #214246 under the version invariant from PR #219687. Supply OpenCL 2.0 in classic CodeGen and CIRGen without changing HIP language options or enabling OpenCL-only Sema restrictions. Assisted-by: Codex / GPT-6 --- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 6 ++++-- clang/lib/CodeGen/CodeGenModule.cpp | 4 +++- clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip | 2 ++ clang/test/CodeGenHIP/hipspv-kernel.cpp | 8 ++++++++ clang/test/SemaHIP/atomic-init.hip | 8 ++++++++ 5 files changed, 25 insertions(+), 3 deletions(-) create mode 100644 clang/test/SemaHIP/atomic-init.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 322b0327fa3abd..6da830289568e6 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -137,8 +137,10 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, cir::CIRDialect::getSourceLanguageAttrName(), cir::SourceLanguageAttr::get(&mlirContext, *sourceLanguage)); if (langOpts.OpenCL || (langOpts.CUDAIsDevice && getTriple().isSPIRV())) { - setOpenCLVersionAttr(cir::CIRDialect::getOpenCLVersionAttrName(), - langOpts.getOpenCLCompatibleVersion()); + // 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); diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index f3d1524ce94e66..40feb4b30f98c1 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 039ab35f1c9062..45483ab5790cb7 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/CodeGenHIP/hipspv-kernel.cpp b/clang/test/CodeGenHIP/hipspv-kernel.cpp index b9a6df12c79688..1ef8a5b19ccfb1 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 00000000000000..6bf90264174841 --- /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
