llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Arseniy Obolenskiy (aobolensk) <details> <summary>Changes</summary> Match classic CodeGen, which emits it on SPIR but not SPIR-V as SPIR v2.0 requires --- Full diff: https://github.com/llvm/llvm-project/pull/228400.diff 8 Files Affected: - (modified) clang/include/clang/CIR/Dialect/IR/CIRDialect.td (+1) - (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+8) - (modified) clang/lib/CIR/Dialect/IR/CIRDialect.cpp (+2-1) - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOpenCLMetadata.cpp (+1) - (modified) clang/test/CIR/CodeGenOpenCL/version.cl (+14) - (modified) clang/test/CIR/IR/invalid-version.cir (+15) - (modified) clang/test/CIR/IR/version.cir (+5) - (modified) clang/test/CIR/Lowering/opencl-version-metadata.cir (+3) ``````````diff diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index d4ec063d8db0f..595ff8545bd7b 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 8bb518f91e933..42275b1bdc0d2 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 fd6cbcc717b43..11f734694c8a6 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 1daee1f17ece2..595ceb90af258 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 abcccae4abe6a..8e0f639a3e4bc 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 f75c4c6365307..1c47e6267f710 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 47603a8dff658..1c940cf1f4d88 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 822e2e48f2735..719a6b478158d 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>>] `````````` </details> https://github.com/llvm/llvm-project/pull/228400 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
