[llvm-branch-commits] [clang] [CIR][OpenCL] Lower OpenCL language version metadata to LLVM dialect (PR #219689)
Akimasa Watanuki via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 3 19:08:20 PDT 2026
https://github.com/Men-cotton updated https://github.com/llvm/llvm-project/pull/219689
>From a4afb736606214830fd77ac69c0714957675227b Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Mon, 8 Jun 2026 20:57:29 +0900
Subject: [PATCH] [CIR][OpenCL] Lower OpenCL language version metadata to LLVM
dialect
Propagate CIR OpenCL language version module attributes as LLVM dialect named metadata before LLVM IR translation.
Assisted-by: Codex / GPT-5.6 Sol
---
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 2 +
.../LowerToLLVMOpenCLMetadata.cpp | 43 +++++++++++++++++++
.../DirectToLLVM/LowerToLLVMOpenCLMetadata.h | 2 +
clang/test/CIR/CodeGenOpenCL/version.cl | 16 +++++++
.../CIR/Lowering/opencl-version-metadata.cir | 16 +++++++
5 files changed, 79 insertions(+)
create mode 100644 clang/test/CIR/Lowering/opencl-version-metadata.cir
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index 6f7509c363fd3..fa81c0a7da9f8 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -4367,6 +4367,8 @@ void ConvertCIRToLLVMPass::processCIRAttrs(mlir::ModuleOp module) {
module->getAttr(cir::CIRDialect::getModuleLevelAsmAttrName()))
module->setAttr(mlir::LLVM::LLVMDialect::getModuleLevelAsmAttrName(),
asmAttr);
+
+ lowerOpenCLModuleMetadataAttrs(module);
}
void ConvertCIRToLLVMPass::runOnOperation() {
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
index e2e7188df009b..1daee1f17ece2 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp
@@ -171,5 +171,48 @@ void OpenCLFunctionMetadataLowering::lower(
convertOpenCLKernelArgMetadata(clArgMetadata, functionMetadata);
}
+static void createOpenCLVersionNamedMetadata(mlir::ModuleOp module,
+ llvm::StringRef name,
+ mlir::LLVM::MDNodeAttr node) {
+ mlir::MLIRContext *ctx = module.getContext();
+ mlir::OpBuilder builder(ctx);
+ builder.setInsertionPointToEnd(module.getBody());
+ mlir::LLVM::NamedMetadataOp::create(builder, module.getLoc(), name,
+ mlir::ArrayAttr::get(ctx, node));
+}
+
+static void emitOpenCLVersionMetadata(mlir::ModuleOp module,
+ llvm::StringRef metadataName,
+ mlir::Attribute versionAttr) {
+ auto version = mlir::cast<cir::OpenCLVersionAttr>(versionAttr);
+ mlir::MLIRContext *ctx = module.getContext();
+ LLVMMetadataNodeBuilder metadataBuilder(ctx);
+
+ unsigned versionMetadata[] = {static_cast<unsigned>(version.getMajor()),
+ static_cast<unsigned>(version.getMinor())};
+ createOpenCLVersionNamedMetadata(module, metadataName,
+ metadataBuilder.getI32Node(versionMetadata));
+}
+
+void lowerOpenCLModuleMetadataAttrs(mlir::ModuleOp module) {
+ struct OpenCLModuleMetadataMapping {
+ llvm::StringRef cirAttrName;
+ llvm::StringRef llvmMetadataName;
+ };
+
+ const OpenCLModuleMetadataMapping moduleMetadataMappings[] = {
+ {cir::CIRDialect::getOpenCLVersionAttrName(), "opencl.ocl.version"},
+ {cir::CIRDialect::getOpenCLCXXVersionAttrName(), "opencl.cxx.version"},
+ };
+
+ for (const OpenCLModuleMetadataMapping &mapping : moduleMetadataMappings) {
+ mlir::Attribute version = module->getAttr(mapping.cirAttrName);
+ if (!version)
+ continue;
+ emitOpenCLVersionMetadata(module, mapping.llvmMetadataName, version);
+ module->removeAttr(mapping.cirAttrName);
+ }
+}
+
} // namespace direct
} // namespace cir
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
index 3865711656c35..b093443224e0e 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h
@@ -32,6 +32,8 @@ class OpenCLFunctionMetadataLowering {
llvm::SmallVector<mlir::Attribute> functionMetadata;
};
+void lowerOpenCLModuleMetadataAttrs(mlir::ModuleOp module);
+
} // namespace direct
} // namespace cir
diff --git a/clang/test/CIR/CodeGenOpenCL/version.cl b/clang/test/CIR/CodeGenOpenCL/version.cl
index 636f52e676604..abcccae4abe6a 100644
--- a/clang/test/CIR/CodeGenOpenCL/version.cl
+++ b/clang/test/CIR/CodeGenOpenCL/version.cl
@@ -2,6 +2,10 @@
// 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
+// RUN: %clang_cc1 -cl-std=CL1.2 -fclangir -emit-llvm -triple spirv64-unknown-unknown %s -o - | FileCheck --check-prefix=CL12-LLVM %s
+// 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
// CL12-CIR: cir.cl.version = #cir.cl.version<1, 2>
// CL30-CIR: cir.cl.version = #cir.cl.version<3, 0>
@@ -9,6 +13,18 @@
// 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>
+// CL12-LLVM-DAG: !opencl.ocl.version = !{[[CL12_VERSION:![0-9]+]]}
+// CL12-LLVM-DAG: [[CL12_VERSION]] = !{i32 1, i32 2}
+// CL30-LLVM-DAG: !opencl.ocl.version = !{[[CL30_VERSION:![0-9]+]]}
+// CL30-LLVM-DAG: [[CL30_VERSION]] = !{i32 3, i32 0}
+// CLCXX10-LLVM-DAG: !opencl.ocl.version = !{[[CLCXX10_VERSION:![0-9]+]]}
+// CLCXX10-LLVM-DAG: !opencl.cxx.version = !{[[CLCXX10_CXX_VERSION:![0-9]+]]}
+// CLCXX10-LLVM-DAG: [[CLCXX10_VERSION]] = !{i32 2, i32 0}
+// CLCXX10-LLVM-DAG: [[CLCXX10_CXX_VERSION]] = !{i32 1, i32 0}
+// CLCXX2021-LLVM-DAG: !opencl.ocl.version = !{[[CLCXX2021_VERSION:![0-9]+]]}
+// 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}
__kernel void version_marker(__global int *out) {
out[0] = 1;
diff --git a/clang/test/CIR/Lowering/opencl-version-metadata.cir b/clang/test/CIR/Lowering/opencl-version-metadata.cir
new file mode 100644
index 0000000000000..822e2e48f2735
--- /dev/null
+++ b/clang/test/CIR/Lowering/opencl-version-metadata.cir
@@ -0,0 +1,16 @@
+// RUN: cir-opt %s -cir-to-llvm -o - | FileCheck %s
+
+module attributes {
+ cir.cl.cxx.version = #cir.cl.version<2021, 0>,
+ cir.cl.version = #cir.cl.version<3, 0>,
+ cir.triple = "x86_64-unknown-linux-gnu"
+} {
+ cir.func @version_marker() {
+ cir.return
+ }
+}
+
+// CHECK-NOT: cir.cl.version
+// CHECK-NOT: cir.cl.cxx.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>>]
More information about the llvm-branch-commits
mailing list