[llvm-branch-commits] [clang] [CIR] Drop lang_address_space's custom parenthesized attribute format (PR #220880)
Henrich Lauko via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 3 06:46:14 PDT 2026
https://github.com/xlauko updated https://github.com/llvm/llvm-project/pull/220880
>From 95f863f3060d8e507e7dc2a13f823b59ed2bb32a Mon Sep 17 00:00:00 2001
From: Henrich Lauko <hlauko at nvidia.com>
Date: Thu, 3 Sep 2026 12:44:31 +0000
Subject: [PATCH] [CIR] Drop lang_address_space's custom parenthesized
attribute format
CIR_LangAddressSpaceAttr overrode its assembly format to
`(` custom<AddressSpaceValue>($value) `)`. The parentheses defeated the
dialect's `#cir.mnemonic<...>` syntax, so the attribute printed as
`#cir<lang_address_space(offload_global)>`. It now uses the bracketed
CIR_EnumAttr default and spells `#cir.lang_address_space<offload_global>`.
The `lang_address_space(x)` spelling inside `!cir.ptr` and `cir.global` is
unaffected, since that goes through MemorySpaceAttrInterface in CIRTypes.cpp.
The attribute-level pair in CIRAttrs.cpp was only reachable from the deleted
format, so it goes away.
---
.../include/clang/CIR/Dialect/IR/CIRAttrs.td | 4 ---
clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 20 -------------
.../kernel-arg-info-single-as.cl | 6 ++--
.../test/CIR/CodeGenOpenCL/kernel-arg-info.cl | 28 +++++++++----------
.../IR/invalid-opencl-kernel-arg-metadata.cir | 12 ++++----
.../CIR/IR/opencl-kernel-arg-metadata.cir | 8 +++---
.../Lowering/opencl-kernel-arg-metadata.cir | 4 +--
7 files changed, 29 insertions(+), 53 deletions(-)
diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
index 126f6a37a739c..f047cf284255c 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td
@@ -1150,10 +1150,6 @@ def CIR_LangAddressSpaceAttr : CIR_EnumAttr<CIR_LangAddressSpace,
}]>
];
- let assemblyFormat = [{
- `(` custom<AddressSpaceValue>($value) `)`
- }];
-
let defaultValue = "cir::LangAddressSpace::Default";
let extraClassDeclaration = [{
diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
index 99da5044751f0..4bb1ce0e02a22 100644
--- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp
@@ -45,26 +45,6 @@ parseFloatLiteral(mlir::AsmParser &parser,
mlir::FailureOr<llvm::APFloat> &value,
cir::FPTypeInterface fpType);
-//===----------------------------------------------------------------------===//
-// AddressSpaceAttr
-//===----------------------------------------------------------------------===//
-
-mlir::ParseResult parseAddressSpaceValue(mlir::AsmParser &p,
- cir::LangAddressSpace &addrSpace) {
- llvm::SMLoc loc = p.getCurrentLocation();
- mlir::FailureOr<cir::LangAddressSpace> result =
- mlir::FieldParser<cir::LangAddressSpace>::parse(p);
- if (mlir::failed(result))
- return p.emitError(loc, "expected address space keyword");
- addrSpace = result.value();
- return mlir::success();
-}
-
-void printAddressSpaceValue(mlir::AsmPrinter &p,
- cir::LangAddressSpace addrSpace) {
- p << cir::stringifyEnum(addrSpace);
-}
-
static mlir::ParseResult parseConstPtr(mlir::AsmParser &parser,
mlir::IntegerAttr &value);
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 c60c36dc61c89..bcdbf8b6e6038 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl
@@ -13,7 +13,7 @@ kernel void spir_addr_space_kernel_args(__global int *G, __constant int *C,
}
// 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)>]
+// 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>]
// LLVM-DAG: define{{.*}} void @spir_addr_space_kernel_args{{.*}} !kernel_arg_addr_space ![[ADDR_SPACES:[0-9]+]]
// LLVM-DAG: ![[ADDR_SPACES]] = !{i32 1, i32 2, i32 3}
@@ -23,7 +23,7 @@ kernel void global_device_host_kernel_args(
__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)>]
+// 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>]
// LLVM-DAG: define{{.*}} void @global_device_host_kernel_args{{.*}} !kernel_arg_addr_space ![[GLOBAL_DEVICE_HOST_ADDR_SPACES:[0-9]+]]
// LLVM-DAG: ![[GLOBAL_DEVICE_HOST_ADDR_SPACES]] = !{i32 5, i32 6}
@@ -35,7 +35,7 @@ kernel void target_address_space_kernel_arg(
// 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)>]
+// CIR-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir.lang_address_space<default>]
// LLVM-DAG: define{{.*}} void @target_address_space_kernel_arg{{.*}}ptr addrspace(5){{.*}} !kernel_arg_addr_space ![[TARGET_ADDR_SPACE:[0-9]+]]
// LLVM-DAG: ![[TARGET_ADDR_SPACE]] = !{i32 0}
diff --git a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
index a00bfec52b122..a3cd52517590b 100644
--- a/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
+++ b/clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl
@@ -28,14 +28,14 @@ kernel void global_qualifier_kernel_args(
// 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: 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: 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*"]
@@ -57,14 +57,14 @@ kernel void constant_kernel_args(constant int *constantintp,
// 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: 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: 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*"]
@@ -90,14 +90,14 @@ kernel void local_qualifier_kernel_args(
// 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: 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: 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*"]
@@ -115,14 +115,14 @@ kernel void private_qualifier_kernel_args(int X, const int constint,
// 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: 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: 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"]
@@ -143,14 +143,14 @@ kernel void typedef_kernel_args(__global unsigned int *X,
// 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: 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: 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*"]
@@ -170,14 +170,14 @@ __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: 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: 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)))*"]
@@ -198,14 +198,14 @@ kernel void signed_char_kernel_args(signed char sc1,
// 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: 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: 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*"]
diff --git a/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir b/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir
index 536187fbbde7d..27c6da3e17fe6 100644
--- a/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir
+++ b/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir
@@ -13,7 +13,7 @@
// expected-error @below {{failed to verify 'access_qual': string array attribute}}
#attr = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<default>],
access_qual = [42 : i32],
type = ["uint*"],
base_type = ["uint*"],
@@ -24,7 +24,7 @@
// expected-error @below {{failed to verify 'type': string array attribute}}
#attr = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<default>],
access_qual = ["none"],
type = [42 : i32],
base_type = ["uint*"],
@@ -35,7 +35,7 @@
// expected-error @below {{failed to verify 'base_type': string array attribute}}
#attr = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<default>],
access_qual = ["none"],
type = ["uint*"],
base_type = [42 : i32],
@@ -46,7 +46,7 @@
// expected-error @below {{failed to verify 'type_qual': string array attribute}}
#attr = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<default>],
access_qual = ["none"],
type = ["uint*"],
base_type = ["uint*"],
@@ -57,7 +57,7 @@
// expected-error @below {{failed to verify 'name': string array attribute or null}}
#attr = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<default>],
access_qual = ["none"],
type = ["uint*"],
base_type = ["uint*"],
@@ -69,7 +69,7 @@
// expected-error @below {{all arrays must have the same number of elements}}
#attr = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<default>],
access_qual = ["none"],
type = ["uint*", "myunsignedint*"],
base_type = ["uint*", "uint*"],
diff --git a/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir b/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir
index 12899dad8f80c..fb6aa92758376 100644
--- a/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir
+++ b/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir
@@ -6,7 +6,7 @@
module {
cir.func @without_names(%arg0: !cir.ptr<!u32i, lang_address_space(offload_global)>, %arg1: !s32i) attributes {
cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<offload_global>, #cir.lang_address_space<default>],
access_qual = ["none", "none"],
type = ["uint*", "int"],
base_type = ["uint*", "int"],
@@ -18,7 +18,7 @@ module {
cir.func @with_names(%data: !cir.ptr<!u32i, lang_address_space(offload_global)>, %count: !s32i) attributes {
cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<
- addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(default)>],
+ addr_space = [#cir.lang_address_space<offload_global>, #cir.lang_address_space<default>],
access_qual = ["none", "none"],
type = ["uint*", "int"],
base_type = ["uint*", "int"],
@@ -46,7 +46,7 @@ module {
// CHECK-SAME: %{{.*}}: !cir.ptr<!u32i, lang_address_space(offload_global)>
// CHECK-SAME: %{{.*}}: !s32i
// CHECK-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata
-// CHECK-SAME: addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(default)>]
+// CHECK-SAME: addr_space = [#cir.lang_address_space<offload_global>, #cir.lang_address_space<default>]
// CHECK-SAME: type = ["uint*", "int"]
// CHECK-SAME: base_type = ["uint*", "int"]
// CHECK-SAME: type_qual = ["restrict", ""]
@@ -56,7 +56,7 @@ module {
// CHECK-SAME: %{{.*}}: !cir.ptr<!u32i, lang_address_space(offload_global)>
// CHECK-SAME: %{{.*}}: !s32i
// CHECK-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata
-// CHECK-SAME: addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(default)>]
+// CHECK-SAME: addr_space = [#cir.lang_address_space<offload_global>, #cir.lang_address_space<default>]
// CHECK-SAME: type = ["uint*", "int"]
// CHECK-SAME: base_type = ["uint*", "int"]
// CHECK-SAME: type_qual = ["restrict", ""]
diff --git a/clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir b/clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir
index 11541fde880cf..541953a6bc606 100644
--- a/clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir
+++ b/clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir
@@ -1,13 +1,13 @@
// RUN: cir-opt %s -cir-to-llvm -o - | FileCheck %s --check-prefix=MLIR
module attributes {cir.triple = "x86_64-unknown-linux-gnu"} {
- cir.func @opencl_kernel_with_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global)>, #cir<lang_address_space(default)>], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""], name = ["data", "count"]>} {
+ cir.func @opencl_kernel_with_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir.lang_address_space<offload_global>, #cir.lang_address_space<default>], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""], name = ["data", "count"]>} {
cir.return
}
cir.func @opencl_target() {
cir.return
}
- cir.func @opencl_alias() alias(@opencl_target) attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir<lang_address_space(offload_global)>], access_qual = ["none"], type = ["uint*"], base_type = ["uint*"], type_qual = [""]>}
+ cir.func @opencl_alias() alias(@opencl_target) attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [#cir.lang_address_space<offload_global>], access_qual = ["none"], type = ["uint*"], base_type = ["uint*"], type_qual = [""]>}
}
// MLIR-LABEL: llvm.func @opencl_kernel_with_names()
More information about the llvm-branch-commits
mailing list