Author: Akimasa Watanuki Date: 2026-08-17T21:41:00+09:00 New Revision: 55e2755aa23cb880865f59618bec6caf97f59d7d
URL: https://github.com/llvm/llvm-project/commit/55e2755aa23cb880865f59618bec6caf97f59d7d DIFF: https://github.com/llvm/llvm-project/commit/55e2755aa23cb880865f59618bec6caf97f59d7d.diff LOG: [CIR][OpenCL] Attach kernel argument metadata to CIR functions (#200581) 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. Added: clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl clang/test/CIR/CodeGenOpenCL/kernel-arg-metadata.cl Modified: clang/lib/CIR/CodeGen/CIRGenFunction.cpp clang/lib/CIR/CodeGen/CIRGenModule.cpp clang/lib/CIR/CodeGen/CIRGenModule.h Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp index 39cf3ee20d52b..b099a4a2e72bf 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp @@ -783,6 +783,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 7d8c54cbfb1d8..98f6ab9755501 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -3328,6 +3328,119 @@ void CIRGenModule::setCIRFunctionAttributesForDefinition( assert(!cir::MissingFeatures::opFuncColdHotAttr()); } +// Maps an AST address space to the OpenCL logical address space kind recorded +// in kernel argument metadata. This mapping is independent of the target +// address space map, allowing consumers to distinguish OpenCL logical address +// spaces even when the target maps them to the same address space. +static cir::LangAddressSpace +getOpenCLKernelArgAddressSpace(LangAS addressSpace) { + switch (addressSpace) { + case LangAS::opencl_global: + return cir::LangAddressSpace::OffloadGlobal; + case LangAS::opencl_constant: + return cir::LangAddressSpace::OffloadConstant; + case LangAS::opencl_local: + return cir::LangAddressSpace::OffloadLocal; + case LangAS::opencl_generic: + return cir::LangAddressSpace::OffloadGeneric; + case LangAS::opencl_global_device: + return cir::LangAddressSpace::OffloadGlobalDevice; + case LangAS::opencl_global_host: + return cir::LangAddressSpace::OffloadGlobalHost; + default: + // All other AST address spaces, including target-specific ones, use the + // OpenCL metadata default, which lowers to SPIR address space ID 0. + return cir::LangAddressSpace::Default; + } +} + +void CIRGenModule::emitOpenCLKernelArgMetadata(cir::FuncOp func, + const clang::FunctionDecl *fd) { + assert(fd && "expected a kernel function declaration"); + const PrintingPolicy &policy = getASTContext().getPrintingPolicy(); + + // Create arrays that represent the kernel argument metadata. Each array has + // one value per kernel argument, in source order. + 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; + }; + + // Type metadata preserves source spelling, while base type metadata uses + // canonical spelling without typedefs. + if (type->isPointerType()) { + QualType pointeeType = type->getPointeeType(); + addressQuals.push_back(cir::LangAddressSpaceAttr::get( + &getMLIRContext(), + getOpenCLKernelArgAddressSpace(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 78612dd2e9de8..b41a2535991a2 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -651,6 +651,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/kernel-arg-info-single-as.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl new file mode 100644 index 0000000000000..4f17d39ccb974 --- /dev/null +++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl @@ -0,0 +1,28 @@ +// Test that OpenCL kernel argument metadata preserves OpenCL logical 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)>] + +// Target-specific address spaces stay on pointer types but do not represent an +// OpenCL language address-space qualifier. +kernel void target_address_space_kernel_arg( + __attribute__((address_space(5))) int *T) {} + +// CIR-LABEL: cir.func{{.*}} @target_address_space_kernel_arg +// CIR-SAME: !cir.ptr<!s32i, target_address_space(5)> +// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(default)>] 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
