[clang] [CIR][OpenCL] Attach kernel argument metadata to CIR functions (PR #200581)
Akimasa Watanuki via cfe-commits
cfe-commits at lists.llvm.org
Mon Jul 27 23:07:52 PDT 2026
https://github.com/Men-cotton updated https://github.com/llvm/llvm-project/pull/200581
>From a3ff15cf9f25af61e6c419ecaad4ca85776838be Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Sun, 24 May 2026 00:56:32 +0900
Subject: [PATCH 1/2] [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 92568dcc3dfd8..e5d3f6fe408c3 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp
@@ -800,6 +800,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 8b9140caa7471..524604b51ba40 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -3266,6 +3266,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 f5ac70a33aef5..cd44709d732fc 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.h
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.h
@@ -657,6 +657,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 = []>
>From b63591dbe14c073c0a3c935d9da77de6420caf06 Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Tue, 28 Jul 2026 14:52:34 +0900
Subject: [PATCH 2/2] fix: map target address spaces to default OpenCL metadata
---
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 30 +++++++++++++------
...ernel-arg-metadata-target-address-space.cl | 5 ----
.../kernel-arg-info-single-as.cl | 9 ++++++
3 files changed, 30 insertions(+), 14 deletions(-)
delete mode 100644 clang/test/CIR/CodeGenOpenCL/invalid-kernel-arg-metadata-target-address-space.cl
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 524604b51ba40..9c3e08e0aa95d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -3266,6 +3266,26 @@ void CIRGenModule::setCIRFunctionAttributesForDefinition(
assert(!cir::MissingFeatures::opFuncColdHotAttr());
}
+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:
+ return cir::LangAddressSpace::Default;
+ }
+}
+
void CIRGenModule::emitOpenCLKernelArgMetadata(cir::FuncOp func,
const clang::FunctionDecl *fd) {
assert(fd && "expected a kernel function declaration");
@@ -3308,17 +3328,9 @@ void CIRGenModule::emitOpenCLKernelArgMetadata(cir::FuncOp func,
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())));
+ getOpenCLKernelArgAddressSpace(pointeeType.getAddressSpace())));
argTypeNames.push_back(
builder.getStringAttr(getTypeSpelling(pointeeType) + "*"));
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
deleted file mode 100644
index f54037175132f..0000000000000
--- a/clang/test/CIR/CodeGenOpenCL/invalid-kernel-arg-metadata-target-address-space.cl
+++ /dev/null
@@ -1,5 +0,0 @@
-// 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
index e18a125098f64..5b560e3727502 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl
@@ -17,3 +17,12 @@ kernel void global_device_host_kernel_args(
// 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)>]
More information about the cfe-commits
mailing list