[llvm-branch-commits] [clang] [CIR][OpenCL] Emit OpenCL language version metadata in CIR (PR #219688)

[Akimasa Watanuki via llvm-branch-commits]

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] [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>