[clang] [CIR][OpenCL] Attach kernel argument metadata to CIR functions (PR #200581)
Akimasa Watanuki via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 12 22:28:12 PDT 2026
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 <mencotton0410 at gmail.com>
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 at +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 = []>
More information about the cfe-commits
mailing list