[clang] [CIR][OpenCL] Emit opencl.spir.version metadata for SPIR targets (PR #228400)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 2 04:28:39 PDT 2026
https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/228400
Match classic CodeGen, which emits it on SPIR but not SPIR-V as SPIR v2.0 requires
>From 9de025e65561ff562979ed0362127a051e1d0376 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Fri, 2 Oct 2026 13:23:45 +0200
Subject: [PATCH] [CIR][OpenCL] Emit opencl.spir.version metadata for SPIR
targets
Match classic CodeGen, which emits it on SPIR but not SPIR-V as SPIR v2.0 requires
---
clang/include/clang/CIR/Dialect/IR/CIRDialect.td | 1 +
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 8 ++++++++
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 3 ++-
.../DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp | 1 +
clang/test/CIR/CodeGenOpenCL/version.cl | 14 ++++++++++++++
clang/test/CIR/IR/invalid-version.cir | 15 +++++++++++++++
clang/test/CIR/IR/version.cir | 5 +++++
.../test/CIR/Lowering/opencl-version-metadata.cir | 3 +++
8 files changed, 49 insertions(+), 1 deletion(-)
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
index d4ec063d8db0fc..595ff8545bd7ba 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td
@@ -122,6 +122,7 @@ def CIR_Dialect : Dialect {
static llvm::StringRef getOpenCLKernelArgMetadataAttrName() { return "cir.cl.kernel_arg_metadata"; }
static llvm::StringRef getOpenCLVersionAttrName() { return "cir.cl.version"; }
static llvm::StringRef getOpenCLCXXVersionAttrName() { return "cir.cl.cxx.version"; }
+ static llvm::StringRef getOpenCLSPIRVersionAttrName() { return "cir.cl.spir.version"; }
static llvm::StringRef getDefaultTlsModelAttrName() { return "cir.default_tls_model"; }
void registerAttributes();
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 8bb518f91e9330..42275b1bdc0d25 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -146,6 +146,14 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext,
if (langOpts.OpenCLCPlusPlus)
setOpenCLVersionAttr(cir::CIRDialect::getOpenCLCXXVersionAttrName(),
langOpts.OpenCLCPlusPlusVersion);
+ // SPIR v2.0 s2.12 requires opencl.spir.version.
+ if (getTriple().isSPIR()) {
+ unsigned spirMajor = version / 100;
+ theModule->setAttr(cir::CIRDialect::getOpenCLSPIRVersionAttrName(),
+ cir::OpenCLVersionAttr::get(&getMLIRContext(),
+ spirMajor,
+ spirMajor > 1 ? 0 : 2));
+ }
}
theModule->setAttr(cir::CIRDialect::getTripleAttrName(),
builder.getStringAttr(getTriple().str()));
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index fd6cbcc717b439..11f734694c8a63 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -114,7 +114,8 @@ Operation *cir::CIRDialect::materializeConstant(mlir::OpBuilder &builder,
static bool isOpenCLVersionAttrName(StringRef attrName) {
return attrName == CIRDialect::getOpenCLVersionAttrName() ||
- attrName == CIRDialect::getOpenCLCXXVersionAttrName();
+ attrName == CIRDialect::getOpenCLCXXVersionAttrName() ||
+ attrName == CIRDialect::getOpenCLSPIRVersionAttrName();
}
static LogicalResult verifyOpenCLVersionAttrPlacement(Operation *op,
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
index 1daee1f17ece2a..595ceb90af258c 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
@@ -203,6 +203,7 @@ void lowerOpenCLModuleMetadataAttrs(mlir::ModuleOp module) {
const OpenCLModuleMetadataMapping moduleMetadataMappings[] = {
{cir::CIRDialect::getOpenCLVersionAttrName(), "opencl.ocl.version"},
{cir::CIRDialect::getOpenCLCXXVersionAttrName(), "opencl.cxx.version"},
+ {cir::CIRDialect::getOpenCLSPIRVersionAttrName(), "opencl.spir.version"},
};
for (const OpenCLModuleMetadataMapping &mapping : moduleMetadataMappings) {
diff --git a/clang/test/CIR/CodeGenOpenCL/version.cl b/clang/test/CIR/CodeGenOpenCL/version.cl
index abcccae4abe6a8..8e0f639a3e4bc0 100644
--- a/clang/test/CIR/CodeGenOpenCL/version.cl
+++ b/clang/test/CIR/CodeGenOpenCL/version.cl
@@ -6,6 +6,12 @@
// RUN: %clang_cc1 -cl-std=CL3.0 -fclangir -emit-llvm -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CL30-LLVM %s
// RUN: %clang_cc1 -x clcpp -cl-std=CLC++ -fclangir -emit-llvm -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CLCXX10-LLVM %s
// RUN: %clang_cc1 -x clcpp -cl-std=CLC++2021 -fclangir -emit-llvm -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CLCXX2021-LLVM %s
+// RUN: %clang_cc1 -cl-std=CL1.2 -fclangir -emit-cir -triple spir-unknown-unknown %s -o - | FileCheck --check-prefix=SPIR12-CIR %s
+// RUN: %clang_cc1 -cl-std=CL3.0 -fclangir -emit-cir -triple spir64-unknown-unknown %s -o - | FileCheck --check-prefix=SPIR30-CIR %s
+// RUN: %clang_cc1 -cl-std=CL1.2 -fclangir -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck --check-prefix=SPIR12-LLVM %s
+// RUN: %clang_cc1 -cl-std=CL3.0 -fclangir -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck --check-prefix=SPIR30-LLVM %s
+// RUN: %clang_cc1 -cl-std=CL1.2 -fclangir -emit-cir -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=SPIRV-CIR %s
+// RUN: %clang_cc1 -cl-std=CL1.2 -fclangir -emit-llvm -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=SPIRV-LLVM %s
// CL12-CIR: cir.cl.version = #cir.cl.version<1, 2>
// CL30-CIR: cir.cl.version = #cir.cl.version<3, 0>
@@ -25,6 +31,14 @@
// CLCXX2021-LLVM-DAG: !opencl.cxx.version = !{[[CLCXX2021_CXX_VERSION:![0-9]+]]}
// CLCXX2021-LLVM-DAG: [[CLCXX2021_VERSION]] = !{i32 3, i32 0}
// CLCXX2021-LLVM-DAG: [[CLCXX2021_CXX_VERSION]] = !{i32 2021, i32 0}
+// SPIR12-CIR: cir.cl.spir.version = #cir.cl.version<1, 2>
+// SPIR30-CIR: cir.cl.spir.version = #cir.cl.version<3, 0>
+// SPIR12-LLVM-DAG: !opencl.spir.version = !{[[SPIR12_VERSION:![0-9]+]]}
+// SPIR12-LLVM-DAG: [[SPIR12_VERSION]] = !{i32 1, i32 2}
+// SPIR30-LLVM-DAG: !opencl.spir.version = !{[[SPIR30_VERSION:![0-9]+]]}
+// SPIR30-LLVM-DAG: [[SPIR30_VERSION]] = !{i32 3, i32 0}
+// SPIRV-CIR-NOT: cir.cl.spir.version
+// SPIRV-LLVM-NOT: !opencl.spir.version
__kernel void version_marker(__global int *out) {
out[0] = 1;
diff --git a/clang/test/CIR/IR/invalid-version.cir b/clang/test/CIR/IR/invalid-version.cir
index f75c4c6365307b..1c47e6267f7103 100644
--- a/clang/test/CIR/IR/invalid-version.cir
+++ b/clang/test/CIR/IR/invalid-version.cir
@@ -71,3 +71,18 @@ module {
// expected-error @below {{cir.cl.cxx.version attribute must be attached to a module}}
cir.func private @not_module_cxx_result() -> (!cir.int<s, 32> {cir.cl.cxx.version = #cir.cl.version<2021, 0>})
}
+
+// -----
+
+// expected-error @below {{expected cir.cl.spir.version to be #cir.cl.version}}
+module attributes {cir.cl.spir.version = "bad"} {
+}
+
+// -----
+
+module {
+ // expected-error @below {{cir.cl.spir.version attribute must be attached to a module}}
+ cir.func @not_module_attr() attributes {cir.cl.spir.version = #cir.cl.version<1, 2>} {
+ cir.return
+ }
+}
diff --git a/clang/test/CIR/IR/version.cir b/clang/test/CIR/IR/version.cir
index 47603a8dff6580..1c940cf1f4d886 100644
--- a/clang/test/CIR/IR/version.cir
+++ b/clang/test/CIR/IR/version.cir
@@ -17,3 +17,8 @@ module attributes {cir.cl.cxx.version = #cir.cl.version<1, 0>, cir.cl.version =
// CHECK: module attributes {cir.cl.cxx.version = #cir.cl.version<2021, 0>, cir.cl.version = #cir.cl.version<3, 0>}
module attributes {cir.cl.cxx.version = #cir.cl.version<2021, 0>, cir.cl.version = #cir.cl.version<3, 0>} { }
+
+// -----
+
+// CHECK: module attributes {cir.cl.spir.version = #cir.cl.version<1, 2>, cir.cl.version = #cir.cl.version<1, 2>}
+module attributes {cir.cl.spir.version = #cir.cl.version<1, 2>, cir.cl.version = #cir.cl.version<1, 2>} { }
diff --git a/clang/test/CIR/Lowering/opencl-version-metadata.cir b/clang/test/CIR/Lowering/opencl-version-metadata.cir
index 822e2e48f2735b..719a6b478158d4 100644
--- a/clang/test/CIR/Lowering/opencl-version-metadata.cir
+++ b/clang/test/CIR/Lowering/opencl-version-metadata.cir
@@ -2,6 +2,7 @@
module attributes {
cir.cl.cxx.version = #cir.cl.version<2021, 0>,
+ cir.cl.spir.version = #cir.cl.version<3, 0>,
cir.cl.version = #cir.cl.version<3, 0>,
cir.triple = "x86_64-unknown-linux-gnu"
} {
@@ -12,5 +13,7 @@ module attributes {
// CHECK-NOT: cir.cl.version
// CHECK-NOT: cir.cl.cxx.version
+// CHECK-NOT: cir.cl.spir.version
// CHECK-DAG: llvm.named_metadata "opencl.ocl.version" [#llvm.md_node<#llvm.md_const<3 : i32>, #llvm.md_const<0 : i32>>]
// CHECK-DAG: llvm.named_metadata "opencl.cxx.version" [#llvm.md_node<#llvm.md_const<2021 : i32>, #llvm.md_const<0 : i32>>]
+// CHECK-DAG: llvm.named_metadata "opencl.spir.version" [#llvm.md_node<#llvm.md_const<3 : i32>, #llvm.md_const<0 : i32>>]
More information about the cfe-commits
mailing list