Author: Akimasa Watanuki Date: 2026-09-12T16:21:54+09:00 New Revision: eff88db8f004bf74c54c65bd39a3fc2fb27969f0
URL: https://github.com/llvm/llvm-project/commit/eff88db8f004bf74c54c65bd39a3fc2fb27969f0 DIFF: https://github.com/llvm/llvm-project/commit/eff88db8f004bf74c54c65bd39a3fc2fb27969f0.diff LOG: [CIR][OpenCL] Lower OpenCL language version metadata to LLVM dialect (#219689) Propagate CIR OpenCL language version module attributes as LLVM dialect named metadata before LLVM IR translation. Assisted-by: Codex / GPT-5.6 Sol Added: clang/test/CIR/Lowering/opencl-version-metadata.cir Modified: clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.h clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip clang/test/CIR/CodeGenOpenCL/version.cl Removed: ################################################################################ diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 5b8cab548f241..186d91599e806 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -4345,6 +4345,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/CodeGenHIP/amdgcnspirv-kernel.hip b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip index 45483ab5790cb..6b4da63db4205 100644 --- a/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip +++ b/clang/test/CIR/CodeGenHIP/amdgcnspirv-kernel.hip @@ -23,3 +23,6 @@ __global__ void kernel_scalar(int a) {} // 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/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>>] _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
