From aa85d46f03fea7962688fd1349ca4a15df6767e6 Mon Sep 17 00:00:00 2001 From: 7mile Date: Sat, 3 Aug 2024 06:17:29 +0800 Subject: [PATCH] [CIR][Dialect] Emit OpenCL kernel argument metadata (#767) Similar to #705, this PR implements the remaining `genKernelArgMetadata()` logic. The attribute `cir.cl.kernel_arg_metadata` is also intentionally placed in the `cir.func`'s `extra_attrs` rather than `cir.func`'s standard `arg_attrs` list. Also, the metadata is stored by `Array` with proper verification on it. See the tablegen doc string for details. This is in order to * keep it side-by-side with `cl.kernel_metadata`. * still emit metadata when kernel has an *empty* arg list (see the test `kernel-arg-meatadata.cl`). * avoid horrors of repeating the long name `cir.cl.kernel_arg_metadata` for `numArgs` times. Because clangir doesn't support OpenCL built-in types and the `half` floating point type yet, their changes and test cases are not included. Corresponding missing feature flag is added. --- .../clang/CIR/Dialect/IR/CIROpenCLAttrs.td | 55 ++++++ clang/include/clang/CIR/MissingFeatures.h | 2 +- clang/lib/CIR/CodeGen/CIRGenFunction.cpp | 3 +- clang/lib/CIR/CodeGen/CIRGenModule.cpp | 171 ++++++++++++++++++ clang/lib/CIR/CodeGen/CIRGenModule.h | 14 ++ clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 37 ++++ .../Lowering/DirectToLLVM/LowerToLLVMIR.cpp | 78 ++++++++ .../OpenCL/kernel-arg-info-single-as.cl | 14 ++ .../CIR/CodeGen/OpenCL/kernel-arg-info.cl | 90 +++++++++ .../CIR/CodeGen/OpenCL/kernel-arg-metadata.cl | 12 ++ clang/test/CIR/IR/invalid.cir | 46 +++++ 11 files changed, 519 insertions(+), 3 deletions(-) create mode 100644 clang/test/CIR/CodeGen/OpenCL/kernel-arg-info-single-as.cl create mode 100644 clang/test/CIR/CodeGen/OpenCL/kernel-arg-info.cl create mode 100644 clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl diff --git a/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td index 1f32701909b7..294f18c9414d 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td @@ -92,4 +92,59 @@ def OpenCLKernelMetadataAttr } +//===----------------------------------------------------------------------===// +// OpenCLKernelArgMetadataAttr +//===----------------------------------------------------------------------===// + +def OpenCLKernelArgMetadataAttr + : CIR_Attr<"OpenCLKernelArgMetadata", "cl.kernel_arg_metadata"> { + + let summary = "OpenCL kernel argument metadata"; + let description = [{ + Provide the required information of an OpenCL kernel argument for the SPIR-V + backend. + + All parameters are arrays, containing the information of the argument in + the same order as they appear in the source code. + + The `addr_space` parameter is an array of I32 that provides the address + space of the argument. It's useful for special types like `image`, which + have implicit global address space. + + Other parameters are arrays of strings that pass through the information + from the source code correspondingly. + + All the fields are mandatory except for `name`, which is optional. + + Example: + ``` + #fn_attr = #cir})> + + cir.func @kernel(%arg0: !s32i) extra(#fn_attr) { + cir.return + } + ``` + }]; + + let parameters = (ins + "ArrayAttr":$addr_space, + "ArrayAttr":$access_qual, + "ArrayAttr":$type, + "ArrayAttr":$base_type, + "ArrayAttr":$type_qual, + OptionalParameter<"ArrayAttr">:$name + ); + + let assemblyFormat = "`<` struct(params) `>`"; + + let genVerifyDecl = 1; +} + #endif // MLIR_CIR_DIALECT_CIR_OPENCL_ATTRS diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index d67989120562..436bc506df7c 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -142,7 +142,7 @@ struct MissingFeatures { static bool getFPFeaturesInEffect() { return false; } static bool cxxABI() { return false; } static bool openCL() { return false; } - static bool openCLGenKernelMetadata() { return false; } + static bool openCLBuiltinTypes() { return false; } static bool CUDA() { return false; } static bool openMP() { return false; } static bool openMPRuntime() { return false; } diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp index 27049934a556..916566936283 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp @@ -1725,8 +1725,7 @@ void CIRGenFunction::buildKernelMetadata(const FunctionDecl *FD, if (!FD->hasAttr() && !FD->hasAttr()) return; - // TODO(cir): CGM.genKernelArgMetadata(Fn, FD, this); - assert(!MissingFeatures::openCLGenKernelMetadata()); + CGM.genKernelArgMetadata(Fn, FD, this); if (!getLangOpts().OpenCL) return; diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index bf2b31cf1510..a8df5255e540 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -3061,3 +3061,174 @@ mlir::cir::SourceLanguage CIRGenModule::getCIRSourceLanguage() { // TODO(cir): support remaining source languages. llvm_unreachable("CIR does not yet support the given source language"); } + +// Returns the address space id that should be produced to the +// kernel_arg_addr_space metadata. This is always fixed to the ids +// as specified in the SPIR 2.0 specification in order to differentiate +// for example in clGetKernelArgInfo() implementation between the address +// spaces with targets without unique mapping to the OpenCL address spaces +// (basically all single AS CPUs). +static unsigned ArgInfoAddressSpace(LangAS AS) { + switch (AS) { + case LangAS::opencl_global: + return 1; + case LangAS::opencl_constant: + return 2; + case LangAS::opencl_local: + return 3; + case LangAS::opencl_generic: + return 4; // Not in SPIR 2.0 specs. + case LangAS::opencl_global_device: + return 5; + case LangAS::opencl_global_host: + return 6; + default: + return 0; // Assume private. + } +} + +void CIRGenModule::genKernelArgMetadata(mlir::cir::FuncOp Fn, + const FunctionDecl *FD, + CIRGenFunction *CGF) { + assert(((FD && CGF) || (!FD && !CGF)) && + "Incorrect use - FD and CGF should either be both null or not!"); + // Create MDNodes that represent the kernel arg metadata. + // Each MDNode is a list in the form of "key", N number of values which is + // the same number of values as their are kernel arguments. + + const PrintingPolicy &Policy = getASTContext().getPrintingPolicy(); + + // Integer values for the kernel argument address space qualifiers. + SmallVector addressQuals; + + // Attrs for the kernel argument access qualifiers (images only). + SmallVector accessQuals; + + // Attrs for the kernel argument type names. + SmallVector argTypeNames; + + // Attrs for the kernel argument base type names. + SmallVector argBaseTypeNames; + + // Attrs for the kernel argument type qualifiers. + SmallVector argTypeQuals; + + // Attrs for the kernel argument names. + SmallVector argNames; + + // OpenCL image and pipe types require special treatments for some metadata + assert(!MissingFeatures::openCLBuiltinTypes()); + + if (FD && CGF) + for (unsigned i = 0, e = FD->getNumParams(); i != e; ++i) { + const ParmVarDecl *parm = FD->getParamDecl(i); + // Get argument name. + argNames.push_back(builder.getStringAttr(parm->getName())); + + if (!getLangOpts().OpenCL) + continue; + QualType ty = parm->getType(); + std::string typeQuals; + + // Get image and pipe access qualifier: + if (ty->isImageType() || ty->isPipeType()) { + llvm_unreachable("NYI"); + } else + accessQuals.push_back(builder.getStringAttr("none")); + + auto getTypeSpelling = [&](QualType Ty) { + auto typeName = Ty.getUnqualifiedType().getAsString(Policy); + + if (Ty.isCanonical()) { + StringRef typeNameRef = typeName; + // Turn "unsigned type" to "utype" + if (typeNameRef.consume_front("unsigned ")) + return std::string("u") + typeNameRef.str(); + if (typeNameRef.consume_front("signed ")) + return typeNameRef.str(); + } + + return typeName; + }; + + if (ty->isPointerType()) { + QualType pointeeTy = ty->getPointeeType(); + + // Get address qualifier. + addressQuals.push_back( + ArgInfoAddressSpace(pointeeTy.getAddressSpace())); + + // Get argument type name. + std::string typeName = getTypeSpelling(pointeeTy) + "*"; + std::string baseTypeName = + getTypeSpelling(pointeeTy.getCanonicalType()) + "*"; + argTypeNames.push_back(builder.getStringAttr(typeName)); + argBaseTypeNames.push_back(builder.getStringAttr(baseTypeName)); + + // Get argument type qualifiers: + if (ty.isRestrictQualified()) + typeQuals = "restrict"; + if (pointeeTy.isConstQualified() || + (pointeeTy.getAddressSpace() == LangAS::opencl_constant)) + typeQuals += typeQuals.empty() ? "const" : " const"; + if (pointeeTy.isVolatileQualified()) + typeQuals += typeQuals.empty() ? "volatile" : " volatile"; + } else { + uint32_t AddrSpc = 0; + bool isPipe = ty->isPipeType(); + if (ty->isImageType() || isPipe) + llvm_unreachable("NYI"); + + addressQuals.push_back(AddrSpc); + + // Get argument type name. + ty = isPipe ? ty->castAs()->getElementType() : ty; + std::string typeName = getTypeSpelling(ty); + std::string baseTypeName = getTypeSpelling(ty.getCanonicalType()); + + // Remove access qualifiers on images + // (as they are inseparable from type in clang implementation, + // but OpenCL spec provides a special query to get access qualifier + // via clGetKernelArgInfo with CL_KERNEL_ARG_ACCESS_QUALIFIER): + if (ty->isImageType()) { + llvm_unreachable("NYI"); + } + + argTypeNames.push_back(builder.getStringAttr(typeName)); + argBaseTypeNames.push_back(builder.getStringAttr(baseTypeName)); + + if (isPipe) + llvm_unreachable("NYI"); + } + argTypeQuals.push_back(builder.getStringAttr(typeQuals)); + } + + bool shouldEmitArgName = getCodeGenOpts().EmitOpenCLArgMetadata || + getCodeGenOpts().HIPSaveKernelArgName; + + if (getLangOpts().OpenCL) { + // The kernel arg name is emitted only when `-cl-kernel-arg-info` is on, + // since it is only used to support `clGetKernelArgInfo` which requires + // `-cl-kernel-arg-info` to work. The other metadata are mandatory because + // they are necessary for OpenCL runtime to set kernel argument. + mlir::ArrayAttr resArgNames = {}; + if (shouldEmitArgName) + resArgNames = builder.getArrayAttr(argNames); + + // Update the function's extra attributes with the kernel argument metadata. + auto value = mlir::cir::OpenCLKernelArgMetadataAttr::get( + Fn.getContext(), builder.getI32ArrayAttr(addressQuals), + builder.getArrayAttr(accessQuals), builder.getArrayAttr(argTypeNames), + builder.getArrayAttr(argBaseTypeNames), + builder.getArrayAttr(argTypeQuals), resArgNames); + mlir::NamedAttrList items{Fn.getExtraAttrs().getElements().getValue()}; + auto oldValue = items.set(value.getMnemonic(), value); + if (oldValue != value) { + Fn.setExtraAttrsAttr(mlir::cir::ExtraFuncAttributesAttr::get( + builder.getContext(), builder.getDictionaryAttr(items))); + } + } else { + if (shouldEmitArgName) + llvm_unreachable("NYI HIPSaveKernelArgName"); + } +} diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index b46befcc949a..fa6da9c9506d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -688,6 +688,20 @@ class CIRGenModule : public CIRGenTypeCache { return *openMPRuntime; } + /// OpenCL v1.2 s5.6.4.6 allows the compiler to store kernel argument + /// information in the program executable. The argument information stored + /// includes the argument name, its type, the address and access qualifiers + /// used. This helper can be used to generate metadata for source code kernel + /// function as well as generated implicitly kernels. If a kernel is generated + /// implicitly null value has to be passed to the last two parameters, + /// otherwise all parameters must have valid non-null values. + /// \param FN is a pointer to IR function being generated. + /// \param FD is a pointer to function declaration if any. + /// \param CGF is a pointer to CIRGenFunction that generates this function. + void genKernelArgMetadata(mlir::cir::FuncOp FN, + const FunctionDecl *FD = nullptr, + CIRGenFunction *CGF = nullptr); + private: // An ordered map of canonical GlobalDecls to their mangled names. llvm::MapVector MangledDeclNames; diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index d9583bba2fd9..f0f5e07aa333 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp @@ -554,6 +554,43 @@ LogicalResult OpenCLKernelMetadataAttr::verify( return success(); } +//===----------------------------------------------------------------------===// +// OpenCLKernelArgMetadataAttr definitions +//===----------------------------------------------------------------------===// + +LogicalResult OpenCLKernelArgMetadataAttr::verify( + ::llvm::function_ref<::mlir::InFlightDiagnostic()> emitError, + ArrayAttr addrSpaces, ArrayAttr accessQuals, ArrayAttr types, + ArrayAttr baseTypes, ArrayAttr typeQuals, ArrayAttr argNames) { + auto isIntArray = [](ArrayAttr elt) { + return llvm::all_of( + elt, [](Attribute elt) { return mlir::isa(elt); }); + }; + auto isStrArray = [](ArrayAttr elt) { + return llvm::all_of( + elt, [](Attribute elt) { return mlir::isa(elt); }); + }; + + if (!isIntArray(addrSpaces)) + return emitError() << "addr_space must be integer arrays"; + if (!llvm::all_of>( + {accessQuals, types, baseTypes, typeQuals}, isStrArray)) + return emitError() + << "access_qual, type, base_type, type_qual must be string arrays"; + if (argNames && !isStrArray(argNames)) { + return emitError() << "name must be a string array"; + } + + if (!llvm::all_of>( + {addrSpaces, accessQuals, types, baseTypes, typeQuals, argNames}, + [&](ArrayAttr arr) { + return !arr || arr.size() == addrSpaces.size(); + })) { + return emitError() << "all arrays must have the same number of elements"; + } + return success(); +} + //===----------------------------------------------------------------------===// // AddressSpaceAttr definitions //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp index 65fe667f6ff9..08aeb902b78e 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMIR.cpp @@ -88,6 +88,11 @@ class CIRDialectLLVMIRTranslationInterface attr.getValue())) { emitOpenCLKernelMetadata(clKernelMetadata, llvmFunc, moduleTranslation); + } else if (auto clArgMetadata = + mlir::dyn_cast( + attr.getValue())) { + emitOpenCLKernelArgMetadata(clArgMetadata, func.getNumArguments(), + llvmFunc, moduleTranslation); } } } @@ -148,6 +153,79 @@ class CIRDialectLLVMIRTranslationInterface llvm::MDNode::get(vmCtx, attrMDArgs)); } } + + void emitOpenCLKernelArgMetadata( + mlir::cir::OpenCLKernelArgMetadataAttr clArgMetadata, unsigned numArgs, + llvm::Function *llvmFunc, + mlir::LLVM::ModuleTranslation &moduleTranslation) const { + auto &vmCtx = moduleTranslation.getLLVMContext(); + + // MDNode for the kernel argument address space qualifiers. + SmallVector addressQuals; + + // MDNode for the kernel argument access qualifiers (images only). + SmallVector accessQuals; + + // MDNode for the kernel argument type names. + SmallVector argTypeNames; + + // MDNode for the kernel argument base type names. + SmallVector argBaseTypeNames; + + // MDNode for the kernel argument type qualifiers. + SmallVector argTypeQuals; + + // MDNode for the kernel argument names. + SmallVector argNames; + + auto lowerStringAttr = [&](mlir::Attribute strAttr) { + return llvm::MDString::get( + vmCtx, mlir::cast(strAttr).getValue()); + }; + + bool shouldEmitArgName = !!clArgMetadata.getName(); + + auto addressSpaceValues = + clArgMetadata.getAddrSpace().getAsValueRange(); + + for (auto &&[i, addrSpace] : llvm::enumerate(addressSpaceValues)) { + // Address space qualifier. + addressQuals.push_back( + llvm::ConstantAsMetadata::get(llvm::ConstantInt::get( + llvm::IntegerType::get(vmCtx, 32), addrSpace))); + + // Access qualifier. + accessQuals.push_back(lowerStringAttr(clArgMetadata.getAccessQual()[i])); + + // Type name. + argTypeNames.push_back(lowerStringAttr(clArgMetadata.getType()[i])); + + // Base type name. + argBaseTypeNames.push_back( + lowerStringAttr(clArgMetadata.getBaseType()[i])); + + // Type qualifier. + argTypeQuals.push_back(lowerStringAttr(clArgMetadata.getTypeQual()[i])); + + // Argument name. + if (shouldEmitArgName) + argNames.push_back(lowerStringAttr(clArgMetadata.getName()[i])); + } + + llvmFunc->setMetadata("kernel_arg_addr_space", + llvm::MDNode::get(vmCtx, addressQuals)); + llvmFunc->setMetadata("kernel_arg_access_qual", + llvm::MDNode::get(vmCtx, accessQuals)); + llvmFunc->setMetadata("kernel_arg_type", + llvm::MDNode::get(vmCtx, argTypeNames)); + llvmFunc->setMetadata("kernel_arg_base_type", + llvm::MDNode::get(vmCtx, argBaseTypeNames)); + llvmFunc->setMetadata("kernel_arg_type_qual", + llvm::MDNode::get(vmCtx, argTypeQuals)); + if (shouldEmitArgName) + llvmFunc->setMetadata("kernel_arg_name", + llvm::MDNode::get(vmCtx, argNames)); + } }; void registerCIRDialectTranslation(mlir::DialectRegistry ®istry) { diff --git a/clang/test/CIR/CodeGen/OpenCL/kernel-arg-info-single-as.cl b/clang/test/CIR/CodeGen/OpenCL/kernel-arg-info-single-as.cl new file mode 100644 index 000000000000..b78ee6dddbf7 --- /dev/null +++ b/clang/test/CIR/CodeGen/OpenCL/kernel-arg-info-single-as.cl @@ -0,0 +1,14 @@ +// Test that the kernel argument info always refers to SPIR address spaces, +// even if the target has only one address space like x86_64 does. +// RUN: %clang_cc1 -fclangir %s -cl-std=CL2.0 -emit-cir -o - -triple x86_64-unknown-linux-gnu -o %t.cir +// RUN: FileCheck %s --input-file=%t.cir --check-prefix=CIR + +// RUN: %clang_cc1 -fclangir %s -cl-std=CL2.0 -emit-llvm -o - -triple x86_64-unknown-linux-gnu -o %t.ll +// RUN: FileCheck %s --input-file=%t.ll --check-prefix=LLVM + +kernel void foo(__global int * G, __constant int *C, __local int *L) { + *G = *C + *L; +} +// CIR: cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-DAG: cir.func @foo({{.+}}) extra(#fn_attr[[KERNEL0]]) +// CIR-ARGINFO-DAG: #fn_attr[[KERNEL0:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-DAG: cir.func @foo({{.+}}) extra(#fn_attr[[KERNEL0]]) + +// LLVM-DAG: define{{.*}} void @foo{{.+}} !kernel_arg_addr_space ![[MD11:[0-9]+]] !kernel_arg_access_qual ![[MD12:[0-9]+]] !kernel_arg_type ![[MD13:[0-9]+]] !kernel_arg_base_type ![[MD13]] !kernel_arg_type_qual ![[MD14:[0-9]+]] { +// LLVM-ARGINFO-DAG: define{{.*}} void @foo{{.+}} !kernel_arg_addr_space ![[MD11:[0-9]+]] !kernel_arg_access_qual ![[MD12:[0-9]+]] !kernel_arg_type ![[MD13:[0-9]+]] !kernel_arg_base_type ![[MD13]] !kernel_arg_type_qual ![[MD14:[0-9]+]] !kernel_arg_name ![[MD15:[0-9]+]] { + +// LLVM-DAG: ![[MD11]] = !{i32 1, i32 1, i32 1, i32 1, i32 2, i32 2, i32 1, i32 1, i32 1, i32 1, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3, i32 0, i32 0, i32 0, i32 0} +// LLVM-DAG: ![[MD12]] = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none", !"none"} +// LLVM-DAG: ![[MD13]] = !{!"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int*", !"int", !"int", !"int", !"int"} +// LLVM-DAG: ![[MD14]] = !{!"", !"restrict", !"const", !"restrict const", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"restrict", !"const", !"restrict const", !"const volatile", !"restrict const volatile", !"volatile", !"restrict volatile", !"", !"", !"", !""} +// LLVM-ARGINFO-DAG: ![[MD15]] = !{!"globalintp", !"globalintrestrictp", !"globalconstintp", !"globalconstintrestrictp", !"constantintp", !"constantintrestrictp", !"globalconstvolatileintp", !"globalconstvolatileintrestrictp", !"globalvolatileintp", !"globalvolatileintrestrictp", !"localintp", !"localintrestrictp", !"localconstintp", !"localconstintrestrictp", !"localconstvolatileintp", !"localconstvolatileintrestrictp", !"localvolatileintp", !"localvolatileintrestrictp", !"X", !"constint", !"constvolatileint", !"volatileint"} + +typedef unsigned int myunsignedint; +kernel void foo4(__global unsigned int * X, __global myunsignedint * Y) { +} + +// CIR-DAG: #fn_attr[[KERNEL4:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-DAG: cir.func @foo4({{.+}}) extra(#fn_attr[[KERNEL4]]) +// CIR-ARGINFO-DAG: #fn_attr[[KERNEL4:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-DAG: cir.func @foo4({{.+}}) extra(#fn_attr[[KERNEL4]]) + +// LLVM-DAG: define{{.*}} void @foo4{{.+}} !kernel_arg_addr_space ![[MD41:[0-9]+]] !kernel_arg_access_qual ![[MD42:[0-9]+]] !kernel_arg_type ![[MD43:[0-9]+]] !kernel_arg_base_type ![[MD44:[0-9]+]] !kernel_arg_type_qual ![[MD45:[0-9]+]] { +// LLVM-ARGINFO-DAG: define{{.*}} void @foo4{{.+}} !kernel_arg_addr_space ![[MD41:[0-9]+]] !kernel_arg_access_qual ![[MD42:[0-9]+]] !kernel_arg_type ![[MD43:[0-9]+]] !kernel_arg_base_type ![[MD44:[0-9]+]] !kernel_arg_type_qual ![[MD45:[0-9]+]] !kernel_arg_name ![[MD46:[0-9]+]] { + +// LLVM-DAG: ![[MD41]] = !{i32 1, i32 1} +// LLVM-DAG: ![[MD42]] = !{!"none", !"none"} +// LLVM-DAG: ![[MD43]] = !{!"uint*", !"myunsignedint*"} +// LLVM-DAG: ![[MD44]] = !{!"uint*", !"uint*"} +// LLVM-DAG: ![[MD45]] = !{!"", !""} +// LLVM-ARGINFO-DAG: ![[MD46]] = !{!"X", !"Y"} + +typedef char char16 __attribute__((ext_vector_type(16))); +__kernel void foo6(__global char16 arg[]) {} + +// CIR-DAG: #fn_attr[[KERNEL6:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-DAG: cir.func @foo6({{.+}}) extra(#fn_attr[[KERNEL6]]) +// CIR-ARGINFO-DAG: #fn_attr[[KERNEL6:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-DAG: cir.func @foo6({{.+}}) extra(#fn_attr[[KERNEL6]]) + +// LLVM-DAG: !kernel_arg_type ![[MD61:[0-9]+]] +// LLVM-ARGINFO-DAG: !kernel_arg_name ![[MD62:[0-9]+]] +// LLVM-DAG: ![[MD61]] = !{!"char16*"} +// LLVM-ARGINFO-DAG: ![[MD62]] = !{!"arg"} + +kernel void foo9(signed char sc1, global const signed char* sc2) {} + +// CIR-DAG: #fn_attr[[KERNEL9:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-DAG: cir.func @foo9({{.+}}) extra(#fn_attr[[KERNEL9]]) +// CIR-ARGINFO-DAG: #fn_attr[[KERNEL9:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-DAG: cir.func @foo9({{.+}}) extra(#fn_attr[[KERNEL9]]) + +// LLVM-DAG: define{{.*}} void @foo9{{.+}} !kernel_arg_addr_space ![[SCHAR_AS_QUAL:[0-9]+]] !kernel_arg_access_qual ![[MD42]] !kernel_arg_type ![[SCHAR_TY:[0-9]+]] !kernel_arg_base_type ![[SCHAR_TY]] !kernel_arg_type_qual ![[SCHAR_QUAL:[0-9]+]] { +// LLVM-ARGINFO-DAG: define{{.*}} void @foo9{{.+}} !kernel_arg_addr_space ![[SCHAR_AS_QUAL:[0-9]+]] !kernel_arg_access_qual ![[MD42]] !kernel_arg_type ![[SCHAR_TY:[0-9]+]] !kernel_arg_base_type ![[SCHAR_TY]] !kernel_arg_type_qual ![[SCHAR_QUAL:[0-9]+]] !kernel_arg_name ![[SCHAR_ARG_NAMES:[0-9]+]] { + +// LLVM-DAG: ![[SCHAR_AS_QUAL]] = !{i32 0, i32 1} +// LLVM-DAG: ![[SCHAR_TY]] = !{!"char", !"char*"} +// LLVM-DAG: ![[SCHAR_QUAL]] = !{!"", !"const"} +// LLVM-ARGINFO-DAG: ![[SCHAR_ARG_NAMES]] = !{!"sc1", !"sc2"} diff --git a/clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl b/clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl new file mode 100644 index 000000000000..7961e0e26244 --- /dev/null +++ b/clang/test/CIR/CodeGen/OpenCL/kernel-arg-metadata.cl @@ -0,0 +1,12 @@ +// 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 + +__kernel void kernel_function() {} + +// CIR: #fn_attr[[ATTR:[0-9]*]] = {{.+}}cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata{{.+}} +// CIR: cir.func @kernel_function() extra(#fn_attr[[ATTR]]) + +// LLVM: define {{.*}}void @kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] { +// LLVM: ![[MD]] = !{} diff --git a/clang/test/CIR/IR/invalid.cir b/clang/test/CIR/IR/invalid.cir index 45dda0a39e42..8386a59ba9bd 100644 --- a/clang/test/CIR/IR/invalid.cir +++ b/clang/test/CIR/IR/invalid.cir @@ -1216,3 +1216,49 @@ cir.func @address_space4(%p : !cir.ptr) { // expected- vec_type_hint = !s32i, vec_type_hint_signedness = 0 > + +// ----- + +// expected-error@+1 {{addr_space must be integer arrays}} +#fn_attr = #cir.cl.kernel_arg_metadata< + addr_space = ["none"], + access_qual = ["none"], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""] +> + +// ----- + +// expected-error@+1 {{access_qual, type, base_type, type_qual must be string arrays}} +#fn_attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = [42 : i32], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""] +> + +// ----- + +// expected-error@+1 {{name must be a string array}} +#fn_attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = ["none"], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""], + name = [33 : i32] +> + +// ----- + +// expected-error@+1 {{all arrays must have the same number of elements}} +#fn_attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = ["none"], + type = ["uint*", "myunsignedint*"], + base_type = ["uint*", "uint*"], + type_qual = [""], + name = ["foo"] +>