[llvm-branch-commits] [clang] [CIR][OpenCL] Emit OpenCL language version metadata in CIR (PR #219688)
Akimasa Watanuki via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sat Sep 5 07:21:07 PDT 2026
https://github.com/Men-cotton updated https://github.com/llvm/llvm-project/pull/219688
>From 86545749214a67d2ec43a26e1b42d5e438a55cc2 Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
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 91f22b2fc5ca1..2afbd7925f595 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 cf82906d6e6c0..f0f86f93ff7d3 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>
>From 1986d1cf0e3fcd2e3168534a9d68a7af7ec03b7d Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
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 ++-
.../CIR/CodeGenHIP/amdgcnspirv-kernel.hip | 28 +++++++++++++++++++
clang/test/CodeGenHIP/hipspv-kernel.cpp | 8 ++++++
clang/test/SemaHIP/atomic-init.hip | 8 ++++++
5 files changed, 51 insertions(+), 3 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip
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 2afbd7925f595..856c3c9b79175 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -136,8 +136,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 17b5f0fe4133d..83e66bebe8a14 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2011,7 +2011,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
new file mode 100644
index 0000000000000..6b4da63db4205
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip
@@ -0,0 +1,28 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir
+
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fclangir \
+// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t.ll
+
+// 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))
+
+// CIR: cir.func{{.*}} @_Z13kernel_scalari{{.*}} cc(spir_kernel)
+// LLVM: define spir_kernel void @_Z13kernel_scalari
+__global__ void kernel_scalar(int a) {}
+
+// CIR: cir.func{{.*}} @_Z9device_fni
+// CIR-NOT: cc(spir_kernel)
+// LLVM: define{{.*}} void @_Z9device_fni
+__device__ void device_fn(int a) {}
+
+// LLVM: !opencl.ocl.version = !{[[OCL:![0-9]+]]}
+// LLVM: [[OCL]] = !{i32 2, i32 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;
+}
More information about the llvm-branch-commits
mailing list