https://github.com/Men-cotton updated https://github.com/llvm/llvm-project/pull/200581
>From 1ac5d63467e54f069cdd344082d82c82b77020b1 Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Sun, 24 May 2026 00:49:55 +0900 Subject: [PATCH 1/7] [CIR][OpenCL] Add kernel argument metadata attribute Add a CIR attribute that carries OpenCL kernel argument metadata in source argument order. Verify that each metadata field has the expected element type and that all present arrays describe the same number of arguments. --- .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 1 + .../clang/CIR/Dialect/IR/CIRDialect.td | 1 + .../clang/CIR/Dialect/IR/CIROpenCLAttrs.td | 46 +++++++++++ clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp | 60 ++++++++++++++ clang/lib/CIR/Dialect/IR/CMakeLists.txt | 1 + .../IR/invalid-opencl-kernel-arg-metadata.cir | 78 +++++++++++++++++++ .../CIR/IR/opencl-kernel-arg-metadata.cir | 27 +++++++ 7 files changed, 214 insertions(+) create mode 100644 clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td create mode 100644 clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp create mode 100644 clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir create mode 100644 clang/test/CIR/IR/opencl-kernel-arg-metadata.cir diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index 4032d8219fff3..19a0c25c8b10e 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -1687,6 +1687,7 @@ def CIR_AnnotationArrayAttr : TypedArrayAttrBase<CIR_AnnotationAttr, "array of cir.annotation attributes">; +include "clang/CIR/Dialect/IR/CIROpenCLAttrs.td" include "clang/CIR/Dialect/IR/CIRCUDAAttrs.td" #endif // CLANG_CIR_DIALECT_IR_CIRATTRS_TD diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index aaa7b48262c80..c20af04f97a1a 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td @@ -82,6 +82,7 @@ def CIR_Dialect : Dialect { static llvm::StringRef getAMDGPUCodeObjectVersionAttrName() { return "cir.amdhsa_code_object_version"; } static llvm::StringRef getAMDGPUPrintfKindAttrName() { return "cir.amdgpu_printf_kind"; } + static llvm::StringRef getOpenCLKernelArgMetadataAttrName() { return "cir.cl.kernel_arg_metadata"; } void registerAttributes(); void registerTypes(); diff --git a/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td new file mode 100644 index 0000000000000..c0ec9c7f28f85 --- /dev/null +++ b/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td @@ -0,0 +1,46 @@ +//===- CIROpenCLAttrs.td - CIR dialect attrs for OpenCL ----*- tablegen -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file declares the CIR dialect attributes for OpenCL. +// +//===----------------------------------------------------------------------===// + +#ifndef CLANG_CIR_DIALECT_IR_CIROPENCLATTRS_TD +#define CLANG_CIR_DIALECT_IR_CIROPENCLATTRS_TD + +//===----------------------------------------------------------------------===// +// OpenCLKernelArgMetadataAttr +//===----------------------------------------------------------------------===// + +def CIR_OpenCLKernelArgMetadataAttr + : CIR_Attr<"OpenCLKernelArgMetadata", "cl.kernel_arg_metadata"> { + let summary = "OpenCL kernel argument metadata"; + let description = [{ + Stores the OpenCL kernel argument metadata emitted to LLVM IR as + `kernel_arg_*` metadata. + + All parameters are arrays containing the argument information in source + order. The `name` field is optional and is emitted only when requested by + `-cl-kernel-arg-info`. + }]; + + let parameters = (ins + "::mlir::ArrayAttr":$addr_space, + "::mlir::ArrayAttr":$access_qual, + "::mlir::ArrayAttr":$type, + "::mlir::ArrayAttr":$base_type, + "::mlir::ArrayAttr":$type_qual, + OptionalParameter<"::mlir::ArrayAttr">:$name + ); + + let assemblyFormat = "`<` struct(params) `>`"; + let genVerifyDecl = 1; + let canHaveIllegalCXXABIType = 0; +} + +#endif // CLANG_CIR_DIALECT_IR_CIROPENCLATTRS_TD diff --git a/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp new file mode 100644 index 0000000000000..57692cd4783b8 --- /dev/null +++ b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp @@ -0,0 +1,60 @@ +//===- CIROpenCLAttrs.cpp - OpenCL specific attributes in CIR -------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file defines the OpenCL-specific attrs in the CIR dialect. +// +//===----------------------------------------------------------------------===// + +#include "clang/CIR/Dialect/IR/CIRAttrs.h" + +#include "mlir/IR/Attributes.h" +#include "mlir/IR/Diagnostics.h" +#include "llvm/ADT/STLExtras.h" + +using namespace mlir; +using namespace cir; + +//===----------------------------------------------------------------------===// +// OpenCLKernelArgMetadataAttr definitions +//===----------------------------------------------------------------------===// + +LogicalResult OpenCLKernelArgMetadataAttr::verify( + function_ref<InFlightDiagnostic()> emitError, ArrayAttr addrSpaces, + ArrayAttr accessQuals, ArrayAttr types, ArrayAttr baseTypes, + ArrayAttr typeQuals, ArrayAttr argNames) { + auto isIntArray = [](ArrayAttr attr) { + return llvm::all_of( + attr, [](Attribute elem) { return mlir::isa<IntegerAttr>(elem); }); + }; + auto isStrArray = [](ArrayAttr attr) { + return llvm::all_of( + attr, [](Attribute elem) { return mlir::isa<StringAttr>(elem); }); + }; + + if (!isIntArray(addrSpaces)) + return emitError() << "addr_space must be an integer array"; + if (!isStrArray(accessQuals)) + return emitError() << "access_qual must be a string array"; + if (!isStrArray(types)) + return emitError() << "type must be a string array"; + if (!isStrArray(baseTypes)) + return emitError() << "base_type must be a string array"; + if (!isStrArray(typeQuals)) + return emitError() << "type_qual must be a string array"; + if (argNames && !isStrArray(argNames)) + return emitError() << "name must be a string array"; + + if (!llvm::all_of(ArrayRef<ArrayAttr>{addrSpaces, accessQuals, types, + baseTypes, typeQuals, argNames}, + [&](ArrayAttr attr) { + return !attr || attr.size() == addrSpaces.size(); + })) + return emitError() << "all arrays must have the same number of elements"; + + return success(); +} diff --git a/clang/lib/CIR/Dialect/IR/CMakeLists.txt b/clang/lib/CIR/Dialect/IR/CMakeLists.txt index 98575941035f2..c8205ebeabf6c 100644 --- a/clang/lib/CIR/Dialect/IR/CMakeLists.txt +++ b/clang/lib/CIR/Dialect/IR/CMakeLists.txt @@ -4,6 +4,7 @@ add_clang_library(MLIRCIR CIRMemorySlot.cpp CIRTypes.cpp CIRDataLayout.cpp + CIROpenCLAttrs.cpp DEPENDS MLIRCIROpsIncGen diff --git a/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir b/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir new file mode 100644 index 0000000000000..23c62c09100f5 --- /dev/null +++ b/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir @@ -0,0 +1,78 @@ +// RUN: cir-opt %s -verify-diagnostics -split-input-file + +// expected-error @below {{addr_space must be an integer array}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = ["none"], + access_qual = ["none"], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""] +> + +// ----- + +// expected-error @below {{access_qual must be a string array}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = [42 : i32], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""] +> + +// ----- + +// expected-error @below {{type must be a string array}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = ["none"], + type = [42 : i32], + base_type = ["uint*"], + type_qual = [""] +> + +// ----- + +// expected-error @below {{base_type must be a string array}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = ["none"], + type = ["uint*"], + base_type = [42 : i32], + type_qual = [""] +> + +// ----- + +// expected-error @below {{type_qual must be a string array}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = ["none"], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [42 : i32] +> + +// ----- + +// expected-error @below {{name must be a string array}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = ["none"], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""], + name = [33 : i32] +> + +// ----- + +// expected-error @below {{all arrays must have the same number of elements}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i32], + access_qual = ["none"], + type = ["uint*", "myunsignedint*"], + base_type = ["uint*", "uint*"], + type_qual = [""], + name = ["foo"] +> diff --git a/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir b/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir new file mode 100644 index 0000000000000..89b3f0722d6f6 --- /dev/null +++ b/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir @@ -0,0 +1,27 @@ +// RUN: cir-opt %s --verify-roundtrip | FileCheck %s + +module { + cir.func @without_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [1 : i32, 0 : i32], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""]>} { + cir.return + } + + cir.func @with_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [1 : i32, 0 : i32], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""], name = ["data", "count"]>} { + cir.return + } +} + +// CHECK-LABEL: cir.func @without_names() +// CHECK-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CHECK-SAME: addr_space = [1 : i32, 0 : i32] +// CHECK-SAME: type = ["uint*", "int"] +// CHECK-SAME: base_type = ["uint*", "int"] +// CHECK-SAME: type_qual = ["restrict", ""] +// CHECK-NOT: name = + +// CHECK-LABEL: cir.func @with_names() +// CHECK-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CHECK-SAME: addr_space = [1 : i32, 0 : i32] +// CHECK-SAME: type = ["uint*", "int"] +// CHECK-SAME: base_type = ["uint*", "int"] +// CHECK-SAME: type_qual = ["restrict", ""] +// CHECK-SAME: name = ["data", "count"] >From cadc32fc2f88c16d70ba683c6c5199daad64200b Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Tue, 26 May 2026 22:17:23 +0900 Subject: [PATCH 2/7] fix: Add zero-argument kernel arg metadata test --- clang/test/CIR/IR/opencl-kernel-arg-metadata.cir | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir b/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir index 89b3f0722d6f6..e6982d5353186 100644 --- a/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir +++ b/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir @@ -8,6 +8,10 @@ module { cir.func @with_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [1 : i32, 0 : i32], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""], name = ["data", "count"]>} { cir.return } + + cir.func @no_args() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [], access_qual = [], type = [], base_type = [], type_qual = []>} { + cir.return + } } // CHECK-LABEL: cir.func @without_names() @@ -25,3 +29,11 @@ module { // CHECK-SAME: base_type = ["uint*", "int"] // CHECK-SAME: type_qual = ["restrict", ""] // CHECK-SAME: name = ["data", "count"] + +// CHECK-LABEL: cir.func @no_args() +// CHECK-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata +// CHECK-SAME: addr_space = [] +// CHECK-SAME: access_qual = [] +// CHECK-SAME: type = [] +// CHECK-SAME: base_type = [] +// CHECK-SAME: type_qual = [] >From 06fa1298eb9d9604bf77d8eaff3088df3ab5099f Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Tue, 26 May 2026 22:17:49 +0900 Subject: [PATCH 3/7] fix: Verify kernel arg addr_space values --- clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp | 19 +++++++++++---- .../IR/invalid-opencl-kernel-arg-metadata.cir | 24 ++++++++++++++++++- 2 files changed, 37 insertions(+), 6 deletions(-) diff --git a/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp index 57692cd4783b8..ac8c01ecc1565 100644 --- a/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp @@ -27,17 +27,26 @@ LogicalResult OpenCLKernelArgMetadataAttr::verify( function_ref<InFlightDiagnostic()> emitError, ArrayAttr addrSpaces, ArrayAttr accessQuals, ArrayAttr types, ArrayAttr baseTypes, ArrayAttr typeQuals, ArrayAttr argNames) { - auto isIntArray = [](ArrayAttr attr) { - return llvm::all_of( - attr, [](Attribute elem) { return mlir::isa<IntegerAttr>(elem); }); + auto isInt32Array = [](ArrayAttr attr) { + return llvm::all_of(attr, [](Attribute elem) { + auto intAttr = mlir::dyn_cast<IntegerAttr>(elem); + return intAttr && intAttr.getType().isInteger(32); + }); + }; + auto isNonNegativeIntArray = [](ArrayAttr attr) { + return llvm::all_of(attr, [](Attribute elem) { + return mlir::cast<IntegerAttr>(elem).getValue().isNonNegative(); + }); }; auto isStrArray = [](ArrayAttr attr) { return llvm::all_of( attr, [](Attribute elem) { return mlir::isa<StringAttr>(elem); }); }; - if (!isIntArray(addrSpaces)) - return emitError() << "addr_space must be an integer array"; + if (!isInt32Array(addrSpaces)) + return emitError() << "addr_space must be an i32 integer array"; + if (!isNonNegativeIntArray(addrSpaces)) + return emitError() << "addr_space values must be non-negative"; if (!isStrArray(accessQuals)) return emitError() << "access_qual must be a string array"; if (!isStrArray(types)) 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 23c62c09100f5..18f239d2e0ee5 100644 --- a/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir +++ b/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir @@ -1,6 +1,6 @@ // RUN: cir-opt %s -verify-diagnostics -split-input-file -// expected-error @below {{addr_space must be an integer array}} +// expected-error @below {{addr_space must be an i32 integer array}} #attr = #cir.cl.kernel_arg_metadata< addr_space = ["none"], access_qual = ["none"], @@ -11,6 +11,28 @@ // ----- +// expected-error @below {{addr_space must be an i32 integer array}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [0 : i64], + access_qual = ["none"], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""] +> + +// ----- + +// expected-error @below {{addr_space values must be non-negative}} +#attr = #cir.cl.kernel_arg_metadata< + addr_space = [-1 : i32], + access_qual = ["none"], + type = ["uint*"], + base_type = ["uint*"], + type_qual = [""] +> + +// ----- + // expected-error @below {{access_qual must be a string array}} #attr = #cir.cl.kernel_arg_metadata< addr_space = [0 : i32], >From 6cfd155d4a6f0c38af846d00540f0a06abeea794 Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Tue, 26 May 2026 23:18:45 +0900 Subject: [PATCH 4/7] fix: Use CIR_LangAddressSpace instead of a raw integer --- .../clang/CIR/Dialect/IR/CIREnumAttr.td | 4 ++- clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp | 16 +++------ clang/lib/CIR/Dialect/IR/CIRTypes.cpp | 2 ++ .../TargetLowering/Targets/AMDGPU.cpp | 2 ++ .../TargetLowering/Targets/NVPTX.cpp | 1 + .../TargetLowering/Targets/SPIRV.cpp | 2 ++ .../IR/invalid-opencl-kernel-arg-metadata.cir | 36 ++++--------------- .../CIR/IR/opencl-kernel-arg-metadata.cir | 8 ++--- 8 files changed, 25 insertions(+), 46 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td index 1de6ffdc08d72..cc6f256ddfef4 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td +++ b/clang/include/clang/CIR/Dialect/IR/CIREnumAttr.td @@ -43,7 +43,9 @@ def CIR_LangAddressSpace : CIR_I32EnumAttr< I32EnumAttrCase<"OffloadLocal", 2, "offload_local">, I32EnumAttrCase<"OffloadGlobal", 3, "offload_global">, I32EnumAttrCase<"OffloadConstant", 4, "offload_constant">, - I32EnumAttrCase<"OffloadGeneric", 5, "offload_generic"> + I32EnumAttrCase<"OffloadGeneric", 5, "offload_generic">, + I32EnumAttrCase<"OffloadGlobalDevice", 6, "offload_global_device">, + I32EnumAttrCase<"OffloadGlobalHost", 7, "offload_global_host"> ]> { let description = [{ Enumerates language-specific address spaces used by CIR. These represent diff --git a/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp index ac8c01ecc1565..0560b3494f1f6 100644 --- a/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp @@ -27,15 +27,9 @@ LogicalResult OpenCLKernelArgMetadataAttr::verify( function_ref<InFlightDiagnostic()> emitError, ArrayAttr addrSpaces, ArrayAttr accessQuals, ArrayAttr types, ArrayAttr baseTypes, ArrayAttr typeQuals, ArrayAttr argNames) { - auto isInt32Array = [](ArrayAttr attr) { + auto isLangAddressSpaceArray = [](ArrayAttr attr) { return llvm::all_of(attr, [](Attribute elem) { - auto intAttr = mlir::dyn_cast<IntegerAttr>(elem); - return intAttr && intAttr.getType().isInteger(32); - }); - }; - auto isNonNegativeIntArray = [](ArrayAttr attr) { - return llvm::all_of(attr, [](Attribute elem) { - return mlir::cast<IntegerAttr>(elem).getValue().isNonNegative(); + return mlir::isa<cir::LangAddressSpaceAttr>(elem); }); }; auto isStrArray = [](ArrayAttr attr) { @@ -43,10 +37,8 @@ LogicalResult OpenCLKernelArgMetadataAttr::verify( attr, [](Attribute elem) { return mlir::isa<StringAttr>(elem); }); }; - if (!isInt32Array(addrSpaces)) - return emitError() << "addr_space must be an i32 integer array"; - if (!isNonNegativeIntArray(addrSpaces)) - return emitError() << "addr_space values must be non-negative"; + if (!isLangAddressSpaceArray(addrSpaces)) + return emitError() << "addr_space must be a language address space array"; if (!isStrArray(accessQuals)) return emitError() << "access_qual must be a string array"; if (!isStrArray(types)) diff --git a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp index e2e754479c654..9c2a40e3681aa 100644 --- a/clang/lib/CIR/Dialect/IR/CIRTypes.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRTypes.cpp @@ -1264,7 +1264,9 @@ cir::LangAddressSpace cir::toCIRLangAddressSpace(clang::LangAS langAS) { case LangAS::opencl_generic: return LangAddressSpace::OffloadGeneric; case LangAS::opencl_global_device: + return LangAddressSpace::OffloadGlobalDevice; case LangAS::opencl_global_host: + return LangAddressSpace::OffloadGlobalHost; case LangAS::sycl_global: case LangAS::sycl_global_device: case LangAS::sycl_global_host: diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp index aa396335dc1cb..f4cdc88d60cf4 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/AMDGPU.cpp @@ -25,6 +25,8 @@ constexpr unsigned AMDGPUAddrSpaceMap[] = { llvm::AMDGPUAS::GLOBAL_ADDRESS, // OffloadGlobal llvm::AMDGPUAS::CONSTANT_ADDRESS, // OffloadConstant llvm::AMDGPUAS::FLAT_ADDRESS, // OffloadGeneric + llvm::AMDGPUAS::GLOBAL_ADDRESS, // OffloadGlobalDevice + llvm::AMDGPUAS::GLOBAL_ADDRESS, // OffloadGlobalHost }; class AMDGPUTargetLoweringInfo : public TargetLoweringInfo { diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp index f38d2b8bfa32d..806e3235b6a8e 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/NVPTX.cpp @@ -17,6 +17,7 @@ constexpr unsigned NVPTXAddrSpaceMap[] = { llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, llvm::NVPTXAS::ADDRESS_SPACE_SHARED, llvm::NVPTXAS::ADDRESS_SPACE_GLOBAL, llvm::NVPTXAS::ADDRESS_SPACE_CONST, llvm::NVPTXAS::ADDRESS_SPACE_GENERIC, + llvm::NVPTXAS::ADDRESS_SPACE_GLOBAL, llvm::NVPTXAS::ADDRESS_SPACE_GLOBAL, }; class NVPTXTargetLoweringInfo : public TargetLoweringInfo { diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp index b759acccd1ac6..5367b4c76e2a0 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/Targets/SPIRV.cpp @@ -21,6 +21,8 @@ constexpr unsigned SPIRVAddrSpaceMap[] = { 1, // CrossWorkgroup 2, // UniformConstant 4, // Generic + 5, // GlobalDevice + 6, // GlobalHost }; class SPIRVTargetLoweringInfo : public TargetLoweringInfo { 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 18f239d2e0ee5..f3cd9bb5a2bdf 100644 --- a/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir +++ b/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir @@ -1,6 +1,6 @@ // RUN: cir-opt %s -verify-diagnostics -split-input-file -// expected-error @below {{addr_space must be an i32 integer array}} +// expected-error @below {{addr_space must be a language address space array}} #attr = #cir.cl.kernel_arg_metadata< addr_space = ["none"], access_qual = ["none"], @@ -11,31 +11,9 @@ // ----- -// expected-error @below {{addr_space must be an i32 integer array}} -#attr = #cir.cl.kernel_arg_metadata< - addr_space = [0 : i64], - access_qual = ["none"], - type = ["uint*"], - base_type = ["uint*"], - type_qual = [""] -> - -// ----- - -// expected-error @below {{addr_space values must be non-negative}} -#attr = #cir.cl.kernel_arg_metadata< - addr_space = [-1 : i32], - access_qual = ["none"], - type = ["uint*"], - base_type = ["uint*"], - type_qual = [""] -> - -// ----- - // expected-error @below {{access_qual must be a string array}} #attr = #cir.cl.kernel_arg_metadata< - addr_space = [0 : i32], + addr_space = [#cir<lang_address_space(default)>], access_qual = [42 : i32], type = ["uint*"], base_type = ["uint*"], @@ -46,7 +24,7 @@ // expected-error @below {{type must be a string array}} #attr = #cir.cl.kernel_arg_metadata< - addr_space = [0 : i32], + addr_space = [#cir<lang_address_space(default)>], access_qual = ["none"], type = [42 : i32], base_type = ["uint*"], @@ -57,7 +35,7 @@ // expected-error @below {{base_type must be a string array}} #attr = #cir.cl.kernel_arg_metadata< - addr_space = [0 : i32], + addr_space = [#cir<lang_address_space(default)>], access_qual = ["none"], type = ["uint*"], base_type = [42 : i32], @@ -68,7 +46,7 @@ // expected-error @below {{type_qual must be a string array}} #attr = #cir.cl.kernel_arg_metadata< - addr_space = [0 : i32], + addr_space = [#cir<lang_address_space(default)>], access_qual = ["none"], type = ["uint*"], base_type = ["uint*"], @@ -79,7 +57,7 @@ // expected-error @below {{name must be a string array}} #attr = #cir.cl.kernel_arg_metadata< - addr_space = [0 : i32], + addr_space = [#cir<lang_address_space(default)>], access_qual = ["none"], type = ["uint*"], base_type = ["uint*"], @@ -91,7 +69,7 @@ // expected-error @below {{all arrays must have the same number of elements}} #attr = #cir.cl.kernel_arg_metadata< - addr_space = [0 : i32], + 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 e6982d5353186..811ac2019a925 100644 --- a/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir +++ b/clang/test/CIR/IR/opencl-kernel-arg-metadata.cir @@ -1,11 +1,11 @@ // RUN: cir-opt %s --verify-roundtrip | FileCheck %s module { - cir.func @without_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [1 : i32, 0 : i32], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""]>} { + cir.func @without_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", ""]>} { cir.return } - cir.func @with_names() attributes {cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata<addr_space = [1 : i32, 0 : i32], access_qual = ["none", "none"], type = ["uint*", "int"], base_type = ["uint*", "int"], type_qual = ["restrict", ""], name = ["data", "count"]>} { + cir.func @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 } @@ -16,7 +16,7 @@ module { // CHECK-LABEL: cir.func @without_names() // CHECK-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata -// CHECK-SAME: addr_space = [1 : i32, 0 : i32] +// 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", ""] @@ -24,7 +24,7 @@ module { // CHECK-LABEL: cir.func @with_names() // CHECK-SAME: cir.cl.kernel_arg_metadata = #cir.cl.kernel_arg_metadata -// CHECK-SAME: addr_space = [1 : i32, 0 : i32] +// 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", ""] >From 23b5ed0ca276fe431e1fb731edc5aac0a0dfd4fe Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Tue, 26 May 2026 23:59:41 +0900 Subject: [PATCH 5/7] fix: Update CIR invalid address space diagnostic --- clang/test/CIR/IR/invalid-addrspace.cir | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CIR/IR/invalid-addrspace.cir b/clang/test/CIR/IR/invalid-addrspace.cir index 882199afd6490..8562aa7a090ab 100644 --- a/clang/test/CIR/IR/invalid-addrspace.cir +++ b/clang/test/CIR/IR/invalid-addrspace.cir @@ -46,7 +46,7 @@ cir.func @lang_address_space_empty(%p : !cir.ptr<!u64i, lang_address_space()>) { // ----- !u64i = !cir.int<u, 64> -// expected-error@+1 {{expected one of [default, offload_private, offload_local, offload_global, offload_constant, offload_generic] for language address space kind}} +// expected-error@+1 {{expected one of [default, offload_private, offload_local, offload_global, offload_constant, offload_generic, offload_global_device, offload_global_host] for language address space kind}} cir.func @lang_address_space_invalid(%p : !cir.ptr<!u64i, lang_address_space(foobar)>) { cir.return } >From 3e4bf5439f1fa788a27621e096006048e92b303c Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Wed, 27 May 2026 23:38:33 +0900 Subject: [PATCH 6/7] fix: constrain CIR OpenCL metadata arrays in TableGen --- .../CIR/Dialect/IR/CIRAttrConstraints.td | 31 +++++++++++++++---- .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 8 +++-- .../clang/CIR/Dialect/IR/CIROpenCLAttrs.td | 12 +++---- clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp | 23 -------------- .../IR/invalid-opencl-kernel-arg-metadata.cir | 12 +++---- 5 files changed, 43 insertions(+), 43 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td index 2548d464fb07f..3cb21dd2f4fc7 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrConstraints.td @@ -71,15 +71,34 @@ def CIR_TryHandlerAttr // ArrayAttr constraints //===----------------------------------------------------------------------===// -def CIR_IntArrayAttr : TypedArrayAttrBase<CIR_AnyIntAttr, - "integer array attribute">; - -def CIR_IntOrGlobalViewArrayAttr : TypedArrayAttrBase<CIR_AnyIntOrGlobalViewAttr, - "integer or global view array attribute">{ +class CIR_TypedArrayAttrBase<Attr element, string summary> + : TypedArrayAttrBase<element, summary> { string cppType = "::mlir::ArrayAttr"; } -def CIR_TryHandlerArrayAttr : TypedArrayAttrBase<CIR_TryHandlerAttr, +class CIR_TypedArrayAttrOrNullBase<Attr element, string summary> + : CIR_TypedArrayAttrBase<element, summary> { + let predicate = Or<[ + CPred<"!$_self">, + CIR_TypedArrayAttrBase<element, summary>.predicate + ]>; + string defaultValue = "::mlir::ArrayAttr()"; +} + +def CIR_IntArrayAttr : CIR_TypedArrayAttrBase<CIR_AnyIntAttr, + "integer array attribute">; + +def CIR_IntOrGlobalViewArrayAttr + : CIR_TypedArrayAttrBase<CIR_AnyIntOrGlobalViewAttr, + "integer or global view array attribute">; + +def CIR_StringArrayAttr + : CIR_TypedArrayAttrBase<StrAttr, "string array attribute">; + +def CIR_StringArrayAttrOrNull + : CIR_TypedArrayAttrOrNullBase<StrAttr, "string array attribute or null">; + +def CIR_TryHandlerArrayAttr : CIR_TypedArrayAttrBase<CIR_TryHandlerAttr, "catch all or unwind or global view array attribute">; #endif // CLANG_CIR_DIALECT_IR_CIRATTRCONSTRAINTS_TD diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index 19a0c25c8b10e..0044c5bd1eadb 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -996,6 +996,10 @@ def CIR_LangAddressSpaceAttr : CIR_EnumAttr<CIR_LangAddressSpace, }]; } +def CIR_LangAddressSpaceArrayAttr + : CIR_TypedArrayAttrBase<CIR_LangAddressSpaceAttr, + "language address space array attribute">; + //===----------------------------------------------------------------------===// // TargetAddressSpaceAttr //===----------------------------------------------------------------------===// @@ -1684,8 +1688,8 @@ def CIR_AnnotationAttr : CIR_Attr<"Annotation", "annotation"> { } def CIR_AnnotationArrayAttr - : TypedArrayAttrBase<CIR_AnnotationAttr, - "array of cir.annotation attributes">; + : CIR_TypedArrayAttrBase<CIR_AnnotationAttr, + "array of cir.annotation attributes">; include "clang/CIR/Dialect/IR/CIROpenCLAttrs.td" include "clang/CIR/Dialect/IR/CIRCUDAAttrs.td" diff --git a/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td index c0ec9c7f28f85..94b41da4c925d 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROpenCLAttrs.td @@ -30,12 +30,12 @@ def CIR_OpenCLKernelArgMetadataAttr }]; let parameters = (ins - "::mlir::ArrayAttr":$addr_space, - "::mlir::ArrayAttr":$access_qual, - "::mlir::ArrayAttr":$type, - "::mlir::ArrayAttr":$base_type, - "::mlir::ArrayAttr":$type_qual, - OptionalParameter<"::mlir::ArrayAttr">:$name + CIR_LangAddressSpaceArrayAttr:$addr_space, + CIR_StringArrayAttr:$access_qual, + CIR_StringArrayAttr:$type, + CIR_StringArrayAttr:$base_type, + CIR_StringArrayAttr:$type_qual, + CIR_StringArrayAttrOrNull:$name ); let assemblyFormat = "`<` struct(params) `>`"; diff --git a/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp index 0560b3494f1f6..fac083c3af7a7 100644 --- a/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIROpenCLAttrs.cpp @@ -27,29 +27,6 @@ LogicalResult OpenCLKernelArgMetadataAttr::verify( function_ref<InFlightDiagnostic()> emitError, ArrayAttr addrSpaces, ArrayAttr accessQuals, ArrayAttr types, ArrayAttr baseTypes, ArrayAttr typeQuals, ArrayAttr argNames) { - auto isLangAddressSpaceArray = [](ArrayAttr attr) { - return llvm::all_of(attr, [](Attribute elem) { - return mlir::isa<cir::LangAddressSpaceAttr>(elem); - }); - }; - auto isStrArray = [](ArrayAttr attr) { - return llvm::all_of( - attr, [](Attribute elem) { return mlir::isa<StringAttr>(elem); }); - }; - - if (!isLangAddressSpaceArray(addrSpaces)) - return emitError() << "addr_space must be a language address space array"; - if (!isStrArray(accessQuals)) - return emitError() << "access_qual must be a string array"; - if (!isStrArray(types)) - return emitError() << "type must be a string array"; - if (!isStrArray(baseTypes)) - return emitError() << "base_type must be a string array"; - if (!isStrArray(typeQuals)) - return emitError() << "type_qual must be a string array"; - if (argNames && !isStrArray(argNames)) - return emitError() << "name must be a string array"; - if (!llvm::all_of(ArrayRef<ArrayAttr>{addrSpaces, accessQuals, types, baseTypes, typeQuals, argNames}, [&](ArrayAttr attr) { 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 f3cd9bb5a2bdf..536187fbbde7d 100644 --- a/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir +++ b/clang/test/CIR/IR/invalid-opencl-kernel-arg-metadata.cir @@ -1,6 +1,6 @@ // RUN: cir-opt %s -verify-diagnostics -split-input-file -// expected-error @below {{addr_space must be a language address space array}} +// expected-error @below {{failed to verify 'addr_space': language address space array attribute}} #attr = #cir.cl.kernel_arg_metadata< addr_space = ["none"], access_qual = ["none"], @@ -11,7 +11,7 @@ // ----- -// expected-error @below {{access_qual must be a string array}} +// expected-error @below {{failed to verify 'access_qual': string array attribute}} #attr = #cir.cl.kernel_arg_metadata< addr_space = [#cir<lang_address_space(default)>], access_qual = [42 : i32], @@ -22,7 +22,7 @@ // ----- -// expected-error @below {{type must be a string array}} +// expected-error @below {{failed to verify 'type': string array attribute}} #attr = #cir.cl.kernel_arg_metadata< addr_space = [#cir<lang_address_space(default)>], access_qual = ["none"], @@ -33,7 +33,7 @@ // ----- -// expected-error @below {{base_type must be a string array}} +// expected-error @below {{failed to verify 'base_type': string array attribute}} #attr = #cir.cl.kernel_arg_metadata< addr_space = [#cir<lang_address_space(default)>], access_qual = ["none"], @@ -44,7 +44,7 @@ // ----- -// expected-error @below {{type_qual must be a string array}} +// expected-error @below {{failed to verify 'type_qual': string array attribute}} #attr = #cir.cl.kernel_arg_metadata< addr_space = [#cir<lang_address_space(default)>], access_qual = ["none"], @@ -55,7 +55,7 @@ // ----- -// expected-error @below {{name must be a string array}} +// 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)>], access_qual = ["none"], >From 5f5642dea46d0becbf0dd79c0bdb14cf3fff31af Mon Sep 17 00:00:00 2001 From: mencotton <[email protected]> Date: Sun, 24 May 2026 00:56:32 +0900 Subject: [PATCH 7/7] [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 | 83 ++++++++++ clang/lib/CIR/CodeGen/CIRGenModule.h | 4 + .../kernel-arg-info-single-as.cl | 19 +++ .../test/CIR/CodeGenOpenCL/kernel-arg-info.cl | 152 ++++++++++++++++++ .../CIR/CodeGenOpenCL/kernel-arg-metadata.cl | 12 ++ 6 files changed, 273 insertions(+) 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 4ecb47a864146..6e84c6a8cdac6 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.cpp @@ -806,6 +806,9 @@ cir::FuncOp CIRGenFunction::generateCode(clang::GlobalDecl gd, cir::FuncOp fn, finishFunction(bodyRange.getEnd()); } + if (getLangOpts().OpenCL && funcDecl->hasAttr<DeviceKernelAttr>()) + cgm.emitOpenCLKernelArgMetadata(fn, funcDecl); + eraseEmptyAndUnusedBlocks(fn); return fn; } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index b377f84e8d370..104c736028548 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -3139,6 +3139,89 @@ 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(); + addressQuals.push_back(cir::LangAddressSpaceAttr::get( + &getMLIRContext(), + cir::toCIRLangAddressSpace(pointeeType.getAddressSpace()))); + + argTypeNames.push_back( + builder.getStringAttr(getTypeSpelling(pointeeType) + "*")); + argBaseTypeNames.push_back(builder.getStringAttr( + getTypeSpelling(pointeeType.getCanonicalType()) + "*")); + + if (type.isRestrictQualified()) + typeQuals = "restrict"; + if (pointeeType.isConstQualified() || + pointeeType.getAddressSpace() == LangAS::opencl_constant) + typeQuals += typeQuals.empty() ? "const" : " const"; + if (pointeeType.isVolatileQualified()) + typeQuals += typeQuals.empty() ? "volatile" : " volatile"; + } else { + addressQuals.push_back(cir::LangAddressSpaceAttr::get( + &getMLIRContext(), cir::LangAddressSpace::Default)); + + argTypeNames.push_back(builder.getStringAttr(getTypeSpelling(type))); + argBaseTypeNames.push_back( + builder.getStringAttr(getTypeSpelling(type.getCanonicalType()))); + } + + argTypeQuals.push_back(builder.getStringAttr(typeQuals)); + } + + mlir::ArrayAttr names; + if (getCodeGenOpts().EmitOpenCLArgMetadata) + names = builder.getArrayAttr(argNames); + + mlir::Attribute metadata = cir::OpenCLKernelArgMetadataAttr::get( + func.getContext(), builder.getArrayAttr(addressQuals), + builder.getArrayAttr(accessQuals), builder.getArrayAttr(argTypeNames), + builder.getArrayAttr(argBaseTypeNames), + builder.getArrayAttr(argTypeQuals), names); + func->setAttr(cir::CIRDialect::getOpenCLKernelArgMetadataAttrName(), + metadata); +} + cir::FuncOp CIRGenModule::getOrCreateCIRFunction( StringRef mangledName, mlir::Type funcType, GlobalDecl gd, bool forVTable, bool dontDefer, bool isThunk, ForDefinition_t isForDefinition, diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.h b/clang/lib/CIR/CodeGen/CIRGenModule.h index fa166c1f39b69..42b3cc55d5786 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.h +++ b/clang/lib/CIR/CodeGen/CIRGenModule.h @@ -658,6 +658,10 @@ class CIRGenModule : public CIRGenTypeCache { void setCIRFunctionAttributesForDefinition(const clang::FunctionDecl *fd, cir::FuncOp f); + /// Generate OpenCL kernel argument metadata for a kernel function. + void emitOpenCLKernelArgMetadata(cir::FuncOp func, + const clang::FunctionDecl *fd); + void emitGlobalDefinition(clang::GlobalDecl gd, mlir::Operation *op = nullptr); void emitGlobalFunctionDefinition(clang::GlobalDecl gd, mlir::Operation *op); diff --git a/clang/test/CIR/CodeGenOpenCL/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 = []> _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
