https://github.com/xlauko created https://github.com/llvm/llvm-project/pull/220880
CIR_LangAddressSpaceAttr overrode its assembly format to `(` custom<AddressSpaceValue>($value) `)`, so the standalone attribute printed as `#cir<lang_address_space(offload_global)>`. The parentheses defeated the dialect's `#cir.mnemonic<...>` syntax and forced the generic fallback. It now uses the bracketed CIR_EnumAttr default and prints as `#cir.lang_address_space<offload_global>`. The override existed to make the enum read as `lang_address_space(x)` inside `!cir.ptr` and `cir.global`. Those spellings are unaffected. Both go through parse/printAddressSpaceValue on MemorySpaceAttrInterface in CIRTypes.cpp, which prints the keyword and parentheses itself. The identically named attribute-level pair in CIRAttrs.cpp was only reachable from the deleted format, so it goes away. --- Stacked on #220879. Based on `users/xlauko/cir-enum-02-tls-model`, so the diff shown is this commit alone. >From 17186daa9cf950273ca61bcb7ea9f6c61b72d5aa Mon Sep 17 00:00:00 2001 From: Henrich Lauko <[email protected]> Date: Thu, 3 Sep 2026 11:29:23 +0000 Subject: [PATCH] [CIR] Drop lang_address_space's custom parenthesized attribute format CIR_LangAddressSpaceAttr overrode its assembly format to `(` custom<AddressSpaceValue>($value) `)`, so the standalone attribute printed as `#cir<lang_address_space(offload_global)>`. The parentheses defeated the dialect's `#cir.mnemonic<...>` syntax and forced the generic fallback. It now uses the bracketed CIR_EnumAttr default and prints as `#cir.lang_address_space<offload_global>`. The override existed to make the enum read as `lang_address_space(x)` inside `!cir.ptr` and `cir.global`. Those spellings are unaffected. Both go through parse/printAddressSpaceValue on MemorySpaceAttrInterface in CIRTypes.cpp, which prints the keyword and parentheses itself. The identically named 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() _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
