Author: Henrich Lauko Date: 2026-09-04T14:18:08+02:00 New Revision: dc6b4e11f136a04699daab6bb19680e7ba6ca427
URL: https://github.com/llvm/llvm-project/commit/dc6b4e11f136a04699daab6bb19680e7ba6ca427 DIFF: https://github.com/llvm/llvm-project/commit/dc6b4e11f136a04699daab6bb19680e7ba6ca427.diff LOG: [CIR] Drop lang_address_space's custom parenthesized attribute format (#220880) 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. Added: Modified: clang/include/clang/CIR/Dialect/IR/CIRAttrs.td clang/include/clang/CIR/Dialect/IR/CIROps.td clang/include/clang/CIR/Dialect/IR/CIRTypes.td clang/lib/CIR/Dialect/IR/CIRAttrs.cpp clang/lib/CIR/Dialect/IR/CIRDialect.cpp clang/lib/CIR/Dialect/IR/CIRTypes.cpp clang/test/CIR/CodeGenOpenCL/kernel-arg-info-single-as.cl clang/test/CIR/CodeGenOpenCL/kernel-arg-info.cl clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir clang/test/CIR/IR/opencl-kernel-arg-metadata.cir clang/test/CIR/Lowering/opencl-kernel-arg-metadata.cir Removed: ################################################################################ diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index 31c934b16d898..6a82b5d4b04d1 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -1208,10 +1208,6 @@ def CIR_LangAddressSpaceAttr : CIR_EnumAttr<CIR_LangAddressSpace, }]> ]; - let assemblyFormat = [{ - `(` custom<AddressSpaceValue>($value) `)` - }]; - let defaultValue = "cir::LangAddressSpace::Default"; let extraClassDeclaration = [{ diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 8ccd086a146fb..5c2c948742f08 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -3406,7 +3406,7 @@ def CIR_GlobalOp : CIR_Op<"global", [ (`tls_refs` `=` $tls_refs^)? (`dso_local` $dso_local^)? (`static_local_guard` `` $static_local_guard^)? - (` ` custom<GlobalAddressSpaceValue>($addr_space)^ )? + (` ` custom<GlobalMemorySpace>($addr_space)^ )? $sym_name (`alias` `(` $aliasee^ `)`)? custom<GlobalOpTypeAndInitialValue>($sym_type, $initial_value, diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index b004f3f202987..a690473a0538b 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td @@ -305,7 +305,7 @@ def CIR_PointerType : CIR_Type<"Pointer", "ptr", [ let assemblyFormat = [{ `<` $pointee - ( `,` ` ` custom<AddressSpaceValue>($addrSpace)^ )? + ( `,` ` ` custom<MemorySpace>($addrSpace)^ )? `>` }]; 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/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index ab5791548262e..c3838413984a6 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -282,11 +282,11 @@ static void printOmittedTerminatorRegion(mlir::OpAsmPrinter &printer, } mlir::OptionalParseResult -parseGlobalAddressSpaceValue(mlir::AsmParser &p, - mlir::ptr::MemorySpaceAttrInterface &attr); +parseGlobalMemorySpace(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &attr); -void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp op, - mlir::ptr::MemorySpaceAttrInterface attr); +void printGlobalMemorySpace(mlir::AsmPrinter &printer, cir::GlobalOp op, + mlir::ptr::MemorySpaceAttrInterface attr); //===----------------------------------------------------------------------===// // AllocaOp diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index 897063171e76a..00ccdfd59192c 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -99,19 +99,13 @@ static void printFuncTypeParams(mlir::AsmPrinter &p, // AddressSpace //===----------------------------------------------------------------------===// -mlir::ParseResult -parseAddressSpaceValue(mlir::AsmParser &p, - mlir::ptr::MemorySpaceAttrInterface &attr); +// Spells the address space of `!cir.ptr` and `cir.global` as either +// `lang_address_space(x)` or `target_address_space(n)`. +mlir::ParseResult parseMemorySpace(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &attr); -void printAddressSpaceValue(mlir::AsmPrinter &printer, - mlir::ptr::MemorySpaceAttrInterface attr); - -// Custom parser/printer for the `addrSpace` parameter in `!cir.ptr`. -mlir::ParseResult parseTargetAddressSpace(mlir::AsmParser &p, - cir::TargetAddressSpaceAttr &attr); - -void printTargetAddressSpace(mlir::AsmPrinter &p, - cir::TargetAddressSpaceAttr attr); +void printMemorySpace(mlir::AsmPrinter &printer, + mlir::ptr::MemorySpaceAttrInterface attr); //===----------------------------------------------------------------------===// // Get autogenerated stuff @@ -1568,9 +1562,8 @@ cir::LangAddressSpace cir::toCIRLangAddressSpace(clang::LangAS langAS) { } } -mlir::ParseResult -parseAddressSpaceValue(mlir::AsmParser &p, - mlir::ptr::MemorySpaceAttrInterface &attr) { +mlir::ParseResult parseMemorySpace(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &attr) { llvm::SMLoc loc = p.getCurrentLocation(); @@ -1617,8 +1610,8 @@ parseAddressSpaceValue(mlir::AsmParser &p, return mlir::success(); } -void printAddressSpaceValue(mlir::AsmPrinter &p, - mlir::ptr::MemorySpaceAttrInterface attr) { +void printMemorySpace(mlir::AsmPrinter &p, + mlir::ptr::MemorySpaceAttrInterface attr) { if (!attr) return; @@ -1637,18 +1630,18 @@ void printAddressSpaceValue(mlir::AsmPrinter &p, } mlir::OptionalParseResult -parseGlobalAddressSpaceValue(mlir::AsmParser &p, - mlir::ptr::MemorySpaceAttrInterface &attr) { +parseGlobalMemorySpace(mlir::AsmParser &p, + mlir::ptr::MemorySpaceAttrInterface &attr) { mlir::SMLoc loc = p.getCurrentLocation(); - if (parseAddressSpaceValue(p, attr).failed()) + if (parseMemorySpace(p, attr).failed()) return p.emitError(loc, "failed to parse Address Space Value for GlobalOp"); return mlir::success(); } -void printGlobalAddressSpaceValue(mlir::AsmPrinter &printer, cir::GlobalOp, - mlir::ptr::MemorySpaceAttrInterface attr) { - printAddressSpaceValue(printer, attr); +void printGlobalMemorySpace(mlir::AsmPrinter &printer, cir::GlobalOp, + mlir::ptr::MemorySpaceAttrInterface attr) { + printMemorySpace(printer, attr); } mlir::ptr::MemorySpaceAttrInterface cir::normalizeDefaultAddressSpace( 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() _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
