llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Akimasa Watanuki (Men-cotton) <details> <summary>Changes</summary> Translate CIR OpenCL kernel argument metadata into the LLVM IR kernel_arg_* metadata attached to kernel functions. Preserve optional argument names so -cl-kernel-arg-info controls the LLVM metadata surface through the CIR attribute. --- Full diff: https://github.com/llvm/llvm-project/pull/200582.diff 4 Files Affected: - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp (+74-1) - (modified) clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl (+12) - (modified) clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl (+60) - (modified) clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl (+12) ``````````diff diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index dbcd0aed88056..fd420cf3153bf 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -15,10 +15,14 @@ #include "mlir/IR/DialectRegistry.h" #include "mlir/Target/LLVMIR/LLVMTranslationInterface.h" #include "mlir/Target/LLVMIR/ModuleTranslation.h" +#include "clang/CIR/Dialect/IR/CIRAttrs.h" #include "clang/CIR/Dialect/IR/CIRDialect.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/IR/Constant.h" +#include "llvm/IR/Function.h" #include "llvm/IR/GlobalVariable.h" +#include "llvm/IR/Metadata.h" +#include "llvm/Support/ErrorHandling.h" using namespace llvm; @@ -75,11 +79,80 @@ class CIRDialectLLVMIRTranslationInterface // Strip the "cir." prefix to get the LLVM attribute name. llvm::StringRef llvmAttrName = attrName.substr(strlen("cir.")); - if (auto strAttr = mlir::dyn_cast<mlir::StringAttr>(attribute.getValue())) + if (auto clArgMetadata = mlir::dyn_cast<cir::OpenCLKernelArgMetadataAttr>( + attribute.getValue())) { + emitOpenCLKernelArgMetadata(clArgMetadata, llvmFunc, moduleTranslation); + } else if (auto strAttr = + mlir::dyn_cast<mlir::StringAttr>(attribute.getValue())) { llvmFunc->addFnAttr(llvmAttrName, strAttr.getValue()); + } return mlir::success(); } + void emitOpenCLKernelArgMetadata( + cir::OpenCLKernelArgMetadataAttr clArgMetadata, llvm::Function *llvmFunc, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { + llvm::LLVMContext &llvmContext = moduleTranslation.getLLVMContext(); + + SmallVector<llvm::Metadata *, 8> addressQuals; + + auto getOpenCLArgInfoAddressSpace = [](cir::LangAddressSpace addressSpace) { + switch (addressSpace) { + case cir::LangAddressSpace::Default: + case cir::LangAddressSpace::OffloadPrivate: + return 0u; + case cir::LangAddressSpace::OffloadGlobal: + return 1u; + case cir::LangAddressSpace::OffloadConstant: + return 2u; + case cir::LangAddressSpace::OffloadLocal: + return 3u; + case cir::LangAddressSpace::OffloadGeneric: + return 4u; + case cir::LangAddressSpace::OffloadGlobalDevice: + return 5u; + case cir::LangAddressSpace::OffloadGlobalHost: + return 6u; + } + llvm_unreachable("unknown CIR language address space"); + }; + + for (cir::LangAddressSpaceAttr addressSpace : + clArgMetadata.getAddrSpace().getAsRange<cir::LangAddressSpaceAttr>()) { + addressQuals.push_back( + llvm::ConstantAsMetadata::get(llvm::ConstantInt::get( + llvm::IntegerType::get(llvmContext, 32), + getOpenCLArgInfoAddressSpace(addressSpace.getValue())))); + } + + llvmFunc->setMetadata("kernel_arg_addr_space", + llvm::MDNode::get(llvmContext, addressQuals)); + llvmFunc->setMetadata( + "kernel_arg_access_qual", + getStringArrayMetadataNode(llvmContext, clArgMetadata.getAccessQual())); + llvmFunc->setMetadata( + "kernel_arg_type", + getStringArrayMetadataNode(llvmContext, clArgMetadata.getType())); + llvmFunc->setMetadata( + "kernel_arg_base_type", + getStringArrayMetadataNode(llvmContext, clArgMetadata.getBaseType())); + llvmFunc->setMetadata( + "kernel_arg_type_qual", + getStringArrayMetadataNode(llvmContext, clArgMetadata.getTypeQual())); + if (clArgMetadata.getName()) + llvmFunc->setMetadata( + "kernel_arg_name", + getStringArrayMetadataNode(llvmContext, clArgMetadata.getName())); + } + + llvm::MDNode *getStringArrayMetadataNode(llvm::LLVMContext &llvmContext, + mlir::ArrayAttr attrs) const { + SmallVector<llvm::Metadata *, 8> metadata; + for (mlir::StringAttr attr : attrs.getAsRange<mlir::StringAttr>()) + metadata.push_back(llvm::MDString::get(llvmContext, attr.getValue())); + return llvm::MDNode::get(llvmContext, metadata); + } + // Translate CIR's module attributes to LLVM's module metadata mlir::LogicalResult amendModule(mlir::ModuleOp mod, mlir::NamedAttribute attribute, diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl index e18a125098f64..f4823b61966cf 100644 --- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl +++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl @@ -2,6 +2,10 @@ // even if the target has only one address space like x86_64 does. // RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu -emit-cir -o %t.cir // RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR +// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu -emit-llvm -o %t.ll +// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM +// RUN: %clang_cc1 %s -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu -emit-llvm -o %t.ogcg.ll +// RUN: FileCheck %s --input-file=%t.ogcg.ll --check-prefix=LLVM kernel void spir_addr_space_kernel_args(__global int *G, __constant int *C, __local int *L) { @@ -11,9 +15,17 @@ kernel void spir_addr_space_kernel_args(__global int *G, __constant int *C, // CIR-LABEL: cir.func{{.*}} @spir_addr_space_kernel_args // CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_constant)>, #cir<lang_address_space(offload_local)>] +// LLVM-LABEL: define{{.*}} void @spir_addr_space_kernel_args +// LLVM-SAME: !kernel_arg_addr_space ![[ADDR_SPACES:[0-9]+]] + kernel void global_device_host_kernel_args( __attribute__((opencl_global_device)) int *D, __attribute__((opencl_global_host)) int *H) {} // CIR-LABEL: cir.func{{.*}} @global_device_host_kernel_args // CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global_device)>, #cir<lang_address_space(offload_global_host)>] + +// LLVM-LABEL: define{{.*}} void @global_device_host_kernel_args +// LLVM-SAME: !kernel_arg_addr_space ![[GLOBAL_DEVICE_HOST_ADDR_SPACES:[0-9]+]] +// LLVM-DAG: ![[ADDR_SPACES]] = !{i32 1, i32 2, i32 3} +// LLVM-DAG: ![[GLOBAL_DEVICE_HOST_ADDR_SPACES]] = !{i32 5, i32 6} diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl index 7788195157715..8098843228226 100644 --- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl +++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl @@ -4,6 +4,15 @@ // RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-cir -cl-kernel-arg-info -o %t.arginfo.cir // RUN: FileCheck %s --input-file=%t.arginfo.cir --check-prefix=CIR-ARGINFO +// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -o %t.ll +// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM +// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -cl-kernel-arg-info -o %t.arginfo.ll +// RUN: FileCheck %s --input-file=%t.arginfo.ll --check-prefix=LLVM-ARGINFO +// RUN: %clang_cc1 %s -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -o %t.ogcg.ll +// RUN: FileCheck %s --input-file=%t.ogcg.ll --check-prefix=LLVM +// RUN: %clang_cc1 %s -cl-std=CL2.0 -triple spirv64-unknown-unknown -emit-llvm -cl-kernel-arg-info -o %t.ogcg.arginfo.ll +// RUN: FileCheck %s --input-file=%t.ogcg.arginfo.ll --check-prefix=LLVM-ARGINFO + kernel void global_qualifier_kernel_args( global int *globalintp, global int *restrict globalintrestrictp, global const int *globalconstintp, @@ -29,6 +38,14 @@ kernel void global_qualifier_kernel_args( // CIR-ARGINFO-SAME: type_qual = ["", "restrict", "const", "restrict const", "const volatile", "restrict const volatile", "volatile", "restrict volatile"] // CIR-ARGINFO-SAME: name = ["globalintp", "globalintrestrictp", "globalconstintp", "globalconstintrestrictp", "globalconstvolatileintp", "globalconstvolatileintrestrictp", "globalvolatileintp", "globalvolatileintrestrictp"] +// LLVM-DAG: define{{.*}} void @global_qualifier_kernel_args{{.+}} !kernel_arg_addr_space ![[GLOBAL_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[GLOBAL_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[GLOBAL_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_type_qual ![[GLOBAL_TYPE_QUALS:[0-9]+]] +// LLVM-ARGINFO-DAG: define{{.*}} void @global_qualifier_kernel_args{{.+}} !kernel_arg_addr_space ![[GLOBAL_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[GLOBAL_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[GLOBAL_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_type_qual ![[GLOBAL_TYPE_QUALS:[0-9]+]] !kernel_arg_name ![[GLOBAL_ARG_NAMES:[0-9]+]] +// LLVM-DAG: ![[GLOBAL_ADDR_SPACES]] = !{i32 1, i32 1, i32 1, i32 1, i32 1, i32 1, i32 1, i32 1} +// LLVM-DAG: ![[GLOBAL_ACCESS_QUALS]] = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none"} +// LLVM-DAG: ![[GLOBAL_ARG_TYPES]] = !{!"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*"} +// LLVM-DAG: ![[GLOBAL_TYPE_QUALS]] = !{!"", !"restrict", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile"} +// LLVM-ARGINFO-DAG: ![[GLOBAL_ARG_NAMES]] = !{!"globalintp", !"globalintrestrictp", !"globalconstintp", !"globalconstintrestrictp", !"globalconstvolatileintp", !"globalconstvolatileintrestrictp", !"globalvolatileintp", !"globalvolatileintrestrictp"} + kernel void constant_kernel_args(constant int *constantintp, constant int *restrict constantintrestrictp) {} @@ -48,6 +65,14 @@ kernel void constant_kernel_args(constant int *constantintp, // CIR-ARGINFO-SAME: type_qual = ["const", "restrict const"] // CIR-ARGINFO-SAME: name = ["constantintp", "constantintrestrictp"] +// LLVM-DAG: define{{.*}} void @constant_kernel_args{{.+}} !kernel_arg_addr_space ![[CONSTANT_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[CONSTANT_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[CONSTANT_ARG_TYPES]] !kernel_arg_type_qual ![[CONSTANT_TYPE_QUALS:[0-9]+]] +// LLVM-ARGINFO-DAG: define{{.*}} void @constant_kernel_args{{.+}} !kernel_arg_addr_space ![[CONSTANT_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[CONSTANT_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[CONSTANT_ARG_TYPES]] !kernel_arg_type_qual ![[CONSTANT_TYPE_QUALS:[0-9]+]] !kernel_arg_name ![[CONSTANT_ARG_NAMES:[0-9]+]] +// LLVM-DAG: ![[CONSTANT_ADDR_SPACES]] = !{i32 2, i32 2} +// LLVM-DAG: ![[CONSTANT_ACCESS_QUALS]] = !{!"none", !"none"} +// LLVM-DAG: ![[CONSTANT_ARG_TYPES]] = !{!"int*", !"int*"} +// LLVM-DAG: ![[CONSTANT_TYPE_QUALS]] = !{!"const", !"restrict const"} +// LLVM-ARGINFO-DAG: ![[CONSTANT_ARG_NAMES]] = !{!"constantintp", !"constantintrestrictp"} + kernel void local_qualifier_kernel_args( local int *localintp, local int *restrict localintrestrictp, local const int *localconstintp, @@ -73,6 +98,11 @@ kernel void local_qualifier_kernel_args( // CIR-ARGINFO-SAME: type_qual = ["", "restrict", "const", "restrict const", "const volatile", "restrict const volatile", "volatile", "restrict volatile"] // CIR-ARGINFO-SAME: name = ["localintp", "localintrestrictp", "localconstintp", "localconstintrestrictp", "localconstvolatileintp", "localconstvolatileintrestrictp", "localvolatileintp", "localvolatileintrestrictp"] +// LLVM-DAG: define{{.*}} void @local_qualifier_kernel_args{{.+}} !kernel_arg_addr_space ![[LOCAL_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[GLOBAL_ACCESS_QUALS]] !kernel_arg_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_base_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_type_qual ![[GLOBAL_TYPE_QUALS]] +// LLVM-ARGINFO-DAG: define{{.*}} void @local_qualifier_kernel_args{{.+}} !kernel_arg_addr_space ![[LOCAL_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[GLOBAL_ACCESS_QUALS]] !kernel_arg_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_base_type ![[GLOBAL_ARG_TYPES]] !kernel_arg_type_qual ![[GLOBAL_TYPE_QUALS]] !kernel_arg_name ![[LOCAL_ARG_NAMES:[0-9]+]] +// LLVM-DAG: ![[LOCAL_ADDR_SPACES]] = !{i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3} +// LLVM-ARGINFO-DAG: ![[LOCAL_ARG_NAMES]] = !{!"localintp", !"localintrestrictp", !"localconstintp", !"localconstintrestrictp", !"localconstvolatileintp", !"localconstvolatileintrestrictp", !"localvolatileintp", !"localvolatileintrestrictp"} + kernel void private_qualifier_kernel_args(int X, const int constint, const volatile int constvolatileint, volatile int volatileint) {} @@ -93,6 +123,14 @@ kernel void private_qualifier_kernel_args(int X, const int constint, // CIR-ARGINFO-SAME: type_qual = ["", "", "", ""] // CIR-ARGINFO-SAME: name = ["X", "constint", "constvolatileint", "volatileint"] +// LLVM-DAG: define{{.*}} void @private_qualifier_kernel_args{{.+}} !kernel_arg_addr_space ![[PRIVATE_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[PRIVATE_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[PRIVATE_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[PRIVATE_ARG_TYPES]] !kernel_arg_type_qual ![[PRIVATE_TYPE_QUALS:[0-9]+]] +// LLVM-ARGINFO-DAG: define{{.*}} void @private_qualifier_kernel_args{{.+}} !kernel_arg_addr_space ![[PRIVATE_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[PRIVATE_ACCESS_QUALS:[0-9]+]] !kernel_arg_type ![[PRIVATE_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[PRIVATE_ARG_TYPES]] !kernel_arg_type_qual ![[PRIVATE_TYPE_QUALS:[0-9]+]] !kernel_arg_name ![[PRIVATE_ARG_NAMES:[0-9]+]] +// LLVM-DAG: ![[PRIVATE_ADDR_SPACES]] = !{i32 0, i32 0, i32 0, i32 0} +// LLVM-DAG: ![[PRIVATE_ACCESS_QUALS]] = !{!"none", !"none", !"none", !"none"} +// LLVM-DAG: ![[PRIVATE_ARG_TYPES]] = !{!"int", !"int", !"int", !"int"} +// LLVM-DAG: ![[PRIVATE_TYPE_QUALS]] = !{!"", !"", !"", !""} +// LLVM-ARGINFO-DAG: ![[PRIVATE_ARG_NAMES]] = !{!"X", !"constint", !"constvolatileint", !"volatileint"} + typedef unsigned int myunsignedint; kernel void typedef_kernel_args(__global unsigned int *X, __global myunsignedint *Y) {} @@ -113,6 +151,15 @@ kernel void typedef_kernel_args(__global unsigned int *X, // CIR-ARGINFO-SAME: type_qual = ["", ""] // CIR-ARGINFO-SAME: name = ["X", "Y"] +// LLVM-DAG: define{{.*}} void @typedef_kernel_args{{.+}} !kernel_arg_addr_space ![[TYPEDEF_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS]] !kernel_arg_type ![[TYPEDEF_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[TYPEDEF_BASE_TYPES:[0-9]+]] !kernel_arg_type_qual ![[TYPEDEF_TYPE_QUALS:[0-9]+]] +// LLVM-ARGINFO-DAG: define{{.*}} void @typedef_kernel_args{{.+}} !kernel_arg_addr_space ![[TYPEDEF_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS]] !kernel_arg_type ![[TYPEDEF_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[TYPEDEF_BASE_TYPES:[0-9]+]] !kernel_arg_type_qual ![[TYPEDEF_TYPE_QUALS:[0-9]+]] !kernel_arg_name ![[TYPEDEF_ARG_NAMES:[0-9]+]] + +// LLVM-DAG: ![[TYPEDEF_ADDR_SPACES]] = !{i32 1, i32 1} +// LLVM-DAG: ![[TYPEDEF_ARG_TYPES]] = !{!"uint*", !"myunsignedint*"} +// LLVM-DAG: ![[TYPEDEF_BASE_TYPES]] = !{!"uint*", !"uint*"} +// LLVM-DAG: ![[TYPEDEF_TYPE_QUALS]] = !{!"", !""} +// LLVM-ARGINFO-DAG: ![[TYPEDEF_ARG_NAMES]] = !{!"X", !"Y"} + typedef char char16 __attribute__((ext_vector_type(16))); __kernel void vector_typedef_kernel_arg(__global char16 arg[]) {} @@ -132,6 +179,11 @@ __kernel void vector_typedef_kernel_arg(__global char16 arg[]) {} // CIR-ARGINFO-SAME: type_qual = [""] // CIR-ARGINFO-SAME: name = ["arg"] +// LLVM-DAG: define{{.*}} void @vector_typedef_kernel_arg{{.+}} !kernel_arg_type ![[VECTOR_TYPEDEF_ARG_TYPES:[0-9]+]] +// LLVM-ARGINFO-DAG: define{{.*}} void @vector_typedef_kernel_arg{{.+}} !kernel_arg_name ![[VECTOR_TYPEDEF_ARG_NAMES:[0-9]+]] +// LLVM-DAG: ![[VECTOR_TYPEDEF_ARG_TYPES]] = !{!"char16*"} +// LLVM-ARGINFO-DAG: ![[VECTOR_TYPEDEF_ARG_NAMES]] = !{!"arg"} + kernel void signed_char_kernel_args(signed char sc1, global const signed char *sc2) {} @@ -150,3 +202,11 @@ kernel void signed_char_kernel_args(signed char sc1, // CIR-ARGINFO-SAME: base_type = ["char", "char*"] // CIR-ARGINFO-SAME: type_qual = ["", "const"] // CIR-ARGINFO-SAME: name = ["sc1", "sc2"] + +// LLVM-DAG: define{{.*}} void @signed_char_kernel_args{{.+}} !kernel_arg_addr_space ![[SIGNED_CHAR_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS]] !kernel_arg_type ![[SIGNED_CHAR_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[SIGNED_CHAR_ARG_TYPES]] !kernel_arg_type_qual ![[SIGNED_CHAR_TYPE_QUALS:[0-9]+]] +// LLVM-ARGINFO-DAG: define{{.*}} void @signed_char_kernel_args{{.+}} !kernel_arg_addr_space ![[SIGNED_CHAR_ADDR_SPACES:[0-9]+]] !kernel_arg_access_qual ![[CONSTANT_ACCESS_QUALS]] !kernel_arg_type ![[SIGNED_CHAR_ARG_TYPES:[0-9]+]] !kernel_arg_base_type ![[SIGNED_CHAR_ARG_TYPES]] !kernel_arg_type_qual ![[SIGNED_CHAR_TYPE_QUALS:[0-9]+]] !kernel_arg_name ![[SIGNED_CHAR_ARG_NAMES:[0-9]+]] + +// LLVM-DAG: ![[SIGNED_CHAR_ADDR_SPACES]] = !{i32 0, i32 1} +// LLVM-DAG: ![[SIGNED_CHAR_ARG_TYPES]] = !{!"char", !"char*"} +// LLVM-DAG: ![[SIGNED_CHAR_TYPE_QUALS]] = !{!"", !"const"} +// LLVM-ARGINFO-DAG: ![[SIGNED_CHAR_ARG_NAMES]] = !{!"sc1", !"sc2"} diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl index dd90ac27d6ec5..85c300d3b005b 100644 --- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl +++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl @@ -1,7 +1,19 @@ // RUN: %clang_cc1 %s -fclangir -triple spirv64-unknown-unknown -emit-cir -o %t.cir // RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR +// RUN: %clang_cc1 %s -fclangir -triple spirv64-unknown-unknown -emit-llvm -o %t.ll +// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM +// RUN: %clang_cc1 %s -triple spirv64-unknown-unknown -emit-llvm -o %t.ogcg.ll +// RUN: FileCheck %s --input-file=%t.ogcg.ll --check-prefix=LLVM __kernel void kernel_function() {} // CIR-LABEL: cir.func @kernel_function() // CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [], access_qual = [], type = [], base_type = [], type_qual = []> + +// LLVM-LABEL: define spir_kernel void @kernel_function() +// LLVM-SAME: !kernel_arg_addr_space ![[EMPTY_ARG_METADATA:[0-9]+]] +// LLVM-SAME: !kernel_arg_access_qual ![[EMPTY_ARG_METADATA]] +// LLVM-SAME: !kernel_arg_type ![[EMPTY_ARG_METADATA]] +// LLVM-SAME: !kernel_arg_base_type ![[EMPTY_ARG_METADATA]] +// LLVM-SAME: !kernel_arg_type_qual ![[EMPTY_ARG_METADATA]] +// LLVM: ![[EMPTY_ARG_METADATA]] = !{} `````````` </details> https://github.com/llvm/llvm-project/pull/200582 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
