https://github.com/Men-cotton updated https://github.com/llvm/llvm-project/pull/200581
>From 285b6e3eba7011ff389431ed3dc1f1c6402ecadf Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Sun, 24 May 2026 00:56:32 +0900 Subject: [PATCH] [CIR][OpenCL] Attach kernel argument metadata to CIR functions Emit the CIR OpenCL kernel argument metadata attribute for kernel functions. Preserve CIR language address-space kinds until lowering and include argument names only when `-cl-kernel-arg-info` is enabled. --- clang/lib/CIR/CodeGen/CIRGenFunction.cpp | 3 + clang/lib/CIR/CodeGen/CIRGenModule.cpp | 91 +++++++++++ clang/lib/CIR/CodeGen/CIRGenModule.h | 4 + ...ernel-arg-metadata-target-address-space.cl | 5 + .../kernel-arg-info-single-as.cl | 19 +++ .../test/CIR/CodeGenOpenCL/kernel-arg-info.cl | 152 ++++++++++++++++++ .../CIR/CodeGenOpenCL/kernel-arg-metadata.cl | 12 ++ 7 files changed, 286 insertions(+) create mode 100644 clang/test/CIR/CodeGenOpenCL/invalid-kernel-arg-metadata-target-address-space.cl create mode 100644 clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl create mode 100644 clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl create mode 100644 clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp index 4b020c96964a7..aaf2823a1ef5e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp @@ -806,6 +806,9 @@ cir::FuncOp CIRGenFunction::generateCode(clang::GlobalDecl gd, cir::FuncOp fn, finishFunction(bodyRange.getEnd()); } + if (getLangOpts().OpenCL && funcDecl->hasAttr<DeviceKernelAttr>()) + cgm.emitOpenCLKernelArgMetadata(fn, funcDecl); + eraseEmptyAndUnusedBlocks(fn); return fn; } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index b377f84e8d370..9456f5c689b3c 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -3139,6 +3139,97 @@ void CIRGenModule::setCIRFunctionAttributesForDefinition( assert(!cir::MissingFeatures::opFuncColdHotAttr()); } +void CIRGenModule::emitOpenCLKernelArgMetadata(cir::FuncOp func, + const clang::FunctionDecl *fd) { + assert(fd && "expected a kernel function declaration"); + const PrintingPolicy &policy = getASTContext().getPrintingPolicy(); + + SmallVector<mlir::Attribute> addressQuals; + SmallVector<mlir::Attribute> accessQuals; + SmallVector<mlir::Attribute> argTypeNames; + SmallVector<mlir::Attribute> argBaseTypeNames; + SmallVector<mlir::Attribute> argTypeQuals; + SmallVector<mlir::Attribute> argNames; + + for (const ParmVarDecl *param : fd->parameters()) { + argNames.push_back(builder.getStringAttr(param->getName())); + + QualType type = param->getType(); + std::string typeQuals; + + if (type->isImageType() || type->isPipeType()) { + errorNYI(param->getSourceRange(), + "OpenCL kernel argument metadata for image and pipe types"); + return; + } + + accessQuals.push_back(builder.getStringAttr("none")); + + auto getTypeSpelling = [&](QualType paramType) { + std::string typeName = paramType.getUnqualifiedType().getAsString(policy); + + if (paramType.isCanonical()) { + StringRef typeNameRef = typeName; + if (typeNameRef.consume_front("unsigned ")) + return std::string("u") + typeNameRef.str(); + if (typeNameRef.consume_front("signed ")) + return typeNameRef.str(); + } + + return typeName; + }; + + if (type->isPointerType()) { + QualType pointeeType = type->getPointeeType(); + if (clang::isTargetAddressSpace(pointeeType.getAddressSpace())) { + errorNYI(param->getSourceRange(), + "OpenCL kernel argument metadata for target-specific " + "address_space(N) kernel parameters; classic CodeGen " + "currently accepts this case"); + return; + } + + addressQuals.push_back(cir::LangAddressSpaceAttr::get( + &getMLIRContext(), + cir::toCIRLangAddressSpace(pointeeType.getAddressSpace()))); + + argTypeNames.push_back( + builder.getStringAttr(getTypeSpelling(pointeeType) + "*")); + argBaseTypeNames.push_back(builder.getStringAttr( + getTypeSpelling(pointeeType.getCanonicalType()) + "*")); + + if (type.isRestrictQualified()) + typeQuals = "restrict"; + if (pointeeType.isConstQualified() || + pointeeType.getAddressSpace() == LangAS::opencl_constant) + typeQuals += typeQuals.empty() ? "const" : " const"; + if (pointeeType.isVolatileQualified()) + typeQuals += typeQuals.empty() ? "volatile" : " volatile"; + } else { + addressQuals.push_back(cir::LangAddressSpaceAttr::get( + &getMLIRContext(), cir::LangAddressSpace::Default)); + + argTypeNames.push_back(builder.getStringAttr(getTypeSpelling(type))); + argBaseTypeNames.push_back( + builder.getStringAttr(getTypeSpelling(type.getCanonicalType()))); + } + + argTypeQuals.push_back(builder.getStringAttr(typeQuals)); + } + + mlir::ArrayAttr names; + if (getCodeGenOpts().EmitOpenCLArgMetadata) + names = builder.getArrayAttr(argNames); + + mlir::Attribute metadata = cir::OpenCLKernelArgMetadataAttr::get( + func.getContext(), builder.getArrayAttr(addressQuals), + builder.getArrayAttr(accessQuals), builder.getArrayAttr(argTypeNames), + builder.getArrayAttr(argBaseTypeNames), + builder.getArrayAttr(argTypeQuals), names); + func->setAttr(cir::CIRDialect::getOpenCLKernelArgMetadataAttrName(), + metadata); +} + cir::FuncOp CIRGenModule::getOrCreateCIRFunction( StringRef mangledName, mlir::Type funcType, GlobalDecl gd, bool forVTable, bool dontDefer, bool isThunk, ForDefinition_t isForDefinition, diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index fa166c1f39b69..42b3cc55d5786 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -658,6 +658,10 @@ class CIRGenModule : public CIRGenTypeCache { void setCIRFunctionAttributesForDefinition(const clang::FunctionDecl *fd, cir::FuncOp f); + /// Generate OpenCL kernel argument metadata for a kernel function. + void emitOpenCLKernelArgMetadata(cir::FuncOp func, + const clang::FunctionDecl *fd); + void emitGlobalDefinition(clang::GlobalDecl gd, mlir::Operation *op = nullptr); void emitGlobalFunctionDefinition(clang::GlobalDecl gd, mlir::Operation *op); diff --git a/clang/test/CIR/CodeGenOpenCL/invalid-kernel-arg-metadata-target-address-space.cl b/clang/test/CIR/CodeGenOpenCL/invalid-kernel-arg-metadata-target-address-space.cl new file mode 100644 index 0000000000000..f54037175132f --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/invalid-kernel-arg-metadata-target-address-space.cl @@ -0,0 +1,5 @@ +// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple x86_64-unknown-linux-gnu -emit-cir -o - -verify + +kernel void invalid_target_addr_space_kernel_arg( + // expected-error@+1 {{ClangIR code gen Not Yet Implemented: OpenCL kernel argument metadata for target-specific address_space(N) kernel parameters; classic CodeGen currently accepts this case}} + __attribute__((address_space(5))) int *T) {} diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl new file mode 100644 index 0000000000000..e18a125098f64 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl @@ -0,0 +1,19 @@ +// Test that OpenCL kernel argument metadata preserves semantic address spaces +// 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 + +kernel void spir_addr_space_kernel_args(__global int *G, __constant int *C, + __local int *L) { + *G = *C + *L; +} + +// 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)>] + +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)>] diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl new file mode 100644 index 0000000000000..7788195157715 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl @@ -0,0 +1,152 @@ +// See also clang/test/CodeGenOpenCL/kernel-arg-info.cl. +// RUN: %clang_cc1 %s -fclangir -cl-std=CL2.0 -triple spirv64-unknown-unknown -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 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 + +kernel void global_qualifier_kernel_args( + global int *globalintp, global int *restrict globalintrestrictp, + global const int *globalconstintp, + global const int *restrict globalconstintrestrictp, + global const volatile int *globalconstvolatileintp, + global const volatile int *restrict globalconstvolatileintrestrictp, + global volatile int *globalvolatileintp, + global volatile int *restrict globalvolatileintrestrictp) {} + +// CIR-LABEL: cir.func{{.*}} @global_qualifier_kernel_args +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-SAME: addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>] +// CIR-SAME: access_qual = ["none", "none", "none", "none", "none", "none", "none", "none"] +// CIR-SAME: type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// CIR-SAME: base_type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// CIR-SAME: type_qual = ["", "restrict", "const", "restrict const", "const volatile", "restrict const volatile", "volatile", "restrict volatile"] +// CIR-ARGINFO-LABEL: cir.func{{.*}} @global_qualifier_kernel_args +// CIR-ARGINFO-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-SAME: addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>] +// CIR-ARGINFO-SAME: access_qual = ["none", "none", "none", "none", "none", "none", "none", "none"] +// CIR-ARGINFO-SAME: type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// CIR-ARGINFO-SAME: base_type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// 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"] + +kernel void constant_kernel_args(constant int *constantintp, + constant int *restrict constantintrestrictp) {} + +// CIR-LABEL: cir.func{{.*}} @constant_kernel_args +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-SAME: addr_space = [#cir<lang_address_space(offload_constant)>, #cir<lang_address_space(offload_constant)>] +// CIR-SAME: access_qual = ["none", "none"] +// CIR-SAME: type = ["int*", "int*"] +// CIR-SAME: base_type = ["int*", "int*"] +// CIR-SAME: type_qual = ["const", "restrict const"] +// CIR-ARGINFO-LABEL: cir.func{{.*}} @constant_kernel_args +// CIR-ARGINFO-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-SAME: addr_space = [#cir<lang_address_space(offload_constant)>, #cir<lang_address_space(offload_constant)>] +// CIR-ARGINFO-SAME: access_qual = ["none", "none"] +// CIR-ARGINFO-SAME: type = ["int*", "int*"] +// CIR-ARGINFO-SAME: base_type = ["int*", "int*"] +// CIR-ARGINFO-SAME: type_qual = ["const", "restrict const"] +// CIR-ARGINFO-SAME: name = ["constantintp", "constantintrestrictp"] + +kernel void local_qualifier_kernel_args( + local int *localintp, local int *restrict localintrestrictp, + local const int *localconstintp, + local const int *restrict localconstintrestrictp, + local const volatile int *localconstvolatileintp, + local const volatile int *restrict localconstvolatileintrestrictp, + local volatile int *localvolatileintp, + local volatile int *restrict localvolatileintrestrictp) {} + +// CIR-LABEL: cir.func{{.*}} @local_qualifier_kernel_args +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-SAME: addr_space = [#cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>] +// CIR-SAME: access_qual = ["none", "none", "none", "none", "none", "none", "none", "none"] +// CIR-SAME: type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// CIR-SAME: base_type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// CIR-SAME: type_qual = ["", "restrict", "const", "restrict const", "const volatile", "restrict const volatile", "volatile", "restrict volatile"] +// CIR-ARGINFO-LABEL: cir.func{{.*}} @local_qualifier_kernel_args +// CIR-ARGINFO-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-SAME: addr_space = [#cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>, #cir<lang_address_space(offload_local)>] +// CIR-ARGINFO-SAME: access_qual = ["none", "none", "none", "none", "none", "none", "none", "none"] +// CIR-ARGINFO-SAME: type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// CIR-ARGINFO-SAME: base_type = ["int*", "int*", "int*", "int*", "int*", "int*", "int*", "int*"] +// 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"] + +kernel void private_qualifier_kernel_args(int X, const int constint, + const volatile int constvolatileint, + volatile int volatileint) {} + +// CIR-LABEL: cir.func{{.*}} @private_qualifier_kernel_args +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-SAME: addr_space = [#cir<lang_address_space(default)>, #cir<lang_address_space(default)>, #cir<lang_address_space(default)>, #cir<lang_address_space(default)>] +// CIR-SAME: access_qual = ["none", "none", "none", "none"] +// CIR-SAME: type = ["int", "int", "int", "int"] +// CIR-SAME: base_type = ["int", "int", "int", "int"] +// CIR-SAME: type_qual = ["", "", "", ""] +// CIR-ARGINFO-LABEL: cir.func{{.*}} @private_qualifier_kernel_args +// CIR-ARGINFO-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-SAME: addr_space = [#cir<lang_address_space(default)>, #cir<lang_address_space(default)>, #cir<lang_address_space(default)>, #cir<lang_address_space(default)>] +// CIR-ARGINFO-SAME: access_qual = ["none", "none", "none", "none"] +// CIR-ARGINFO-SAME: type = ["int", "int", "int", "int"] +// CIR-ARGINFO-SAME: base_type = ["int", "int", "int", "int"] +// CIR-ARGINFO-SAME: type_qual = ["", "", "", ""] +// CIR-ARGINFO-SAME: name = ["X", "constint", "constvolatileint", "volatileint"] + +typedef unsigned int myunsignedint; +kernel void typedef_kernel_args(__global unsigned int *X, + __global myunsignedint *Y) {} + +// CIR-LABEL: cir.func{{.*}} @typedef_kernel_args +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-SAME: addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>] +// CIR-SAME: access_qual = ["none", "none"] +// CIR-SAME: type = ["uint*", "myunsignedint*"] +// CIR-SAME: base_type = ["uint*", "uint*"] +// CIR-SAME: type_qual = ["", ""] +// CIR-ARGINFO-LABEL: cir.func{{.*}} @typedef_kernel_args +// CIR-ARGINFO-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-SAME: addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(offload_global)>] +// CIR-ARGINFO-SAME: access_qual = ["none", "none"] +// CIR-ARGINFO-SAME: type = ["uint*", "myunsignedint*"] +// CIR-ARGINFO-SAME: base_type = ["uint*", "uint*"] +// CIR-ARGINFO-SAME: type_qual = ["", ""] +// CIR-ARGINFO-SAME: name = ["X", "Y"] + +typedef char char16 __attribute__((ext_vector_type(16))); +__kernel void vector_typedef_kernel_arg(__global char16 arg[]) {} + +// CIR-LABEL: cir.func{{.*}} @vector_typedef_kernel_arg +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-SAME: addr_space = [#cir<lang_address_space(offload_global)>] +// CIR-SAME: access_qual = ["none"] +// CIR-SAME: type = ["char16*"] +// CIR-SAME: base_type = ["char __attribute__((ext_vector_type(16)))*"] +// CIR-SAME: type_qual = [""] +// CIR-ARGINFO-LABEL: cir.func{{.*}} @vector_typedef_kernel_arg +// CIR-ARGINFO-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-SAME: addr_space = [#cir<lang_address_space(offload_global)>] +// CIR-ARGINFO-SAME: access_qual = ["none"] +// CIR-ARGINFO-SAME: type = ["char16*"] +// CIR-ARGINFO-SAME: base_type = ["char __attribute__((ext_vector_type(16)))*"] +// CIR-ARGINFO-SAME: type_qual = [""] +// CIR-ARGINFO-SAME: name = ["arg"] + +kernel void signed_char_kernel_args(signed char sc1, + global const signed char *sc2) {} + +// CIR-LABEL: cir.func{{.*}} @signed_char_kernel_args +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-SAME: addr_space = [#cir<lang_address_space(default)>, #cir<lang_address_space(offload_global)>] +// CIR-SAME: access_qual = ["none", "none"] +// CIR-SAME: type = ["char", "char*"] +// CIR-SAME: base_type = ["char", "char*"] +// CIR-SAME: type_qual = ["", "const"] +// CIR-ARGINFO-LABEL: cir.func{{.*}} @signed_char_kernel_args +// CIR-ARGINFO-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CIR-ARGINFO-SAME: addr_space = [#cir<lang_address_space(default)>, #cir<lang_address_space(offload_global)>] +// CIR-ARGINFO-SAME: access_qual = ["none", "none"] +// CIR-ARGINFO-SAME: type = ["char", "char*"] +// CIR-ARGINFO-SAME: base_type = ["char", "char*"] +// CIR-ARGINFO-SAME: type_qual = ["", "const"] +// CIR-ARGINFO-SAME: name = ["sc1", "sc2"] diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl new file mode 100644 index 0000000000000..b1ae2d8250b69 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/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 + +extern __kernel void alias_kernel_function(void) + __attribute__((alias("kernel_function"))); + +// CIR-LABEL: cir.func @alias_kernel_function() alias(@kernel_function) + +__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 = []> _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
