https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/229330
Add `!cir.metadata` and `cir.metadata_as_value` so `cir.call_llvm_intrinsic` can take metadata arguments, then use them to lower `__builtin_amdgcn_processor_is` and `__builtin_amdgcn_is_invocable` In AMDGCN-flavored SPIR-V these intrinsics are used to lower spec constants (value will be resolved later on JIT stage for the particular GPU) >From 950bbea47117c1eca8d5eee56f94ee1ba51753ac Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Tue, 6 Oct 2026 09:41:32 +0200 Subject: [PATCH] [CIR][AMDGPU][SPIR-V] Support metadata intrinsic args for feature predicates Add `!cir.metadata` and `cir.metadata_as_value` so `cir.call_llvm_intrinsic` can take metadata arguments, then use them to emit `__builtin_amdgcn_processor_is` and `__builtin_amdgcn_is_invocable` --- .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 39 ++++++++++++++++ clang/include/clang/CIR/Dialect/IR/CIROps.td | 28 +++++++++++- .../include/clang/CIR/Dialect/IR/CIRTypes.td | 12 +++++ clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 36 +++++++++++++++ clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 13 ++++++ .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 22 +++++++++ .../builtin-amdgcn-feature-predicates.c | 45 +++++++++++++++++++ clang/test/CIR/IR/invalid-metadata.cir | 21 +++++++++ .../test/CIR/Lowering/call-llvm-intrinsic.cir | 22 +++++++++ 9 files changed, 236 insertions(+), 2 deletions(-) create mode 100644 clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c create mode 100644 clang/test/CIR/IR/invalid-metadata.cir diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index f6aaac5ab4cab..e7e1f05f86dac 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -1960,6 +1960,45 @@ def CIR_EhUnexpectedAttr : CIR_UnitAttr<"EhUnexpected", "eh_unexpected"> { let canHaveIllegalCXXABIType = 0; } +//===----------------------------------------------------------------------===// +// MDStringAttr & MDNodeAttr +//===----------------------------------------------------------------------===// + +def CIR_MDStringAttr : CIR_Attr<"MDString", "md_string"> { + let summary = "Metadata string"; + let description = [{ + A metadata string, lowered to `llvm::MDString`. + + Example: + ```mlir + #cir.md_string<"agent"> + ``` + }]; + + let parameters = (ins "mlir::StringAttr":$value); + let assemblyFormat = [{ `<` $value `>` }]; + let canHaveIllegalCXXABIType = 0; +} + +def CIR_MDNodeAttr : CIR_Attr<"MDNode", "md_node"> { + let summary = "Metadata node"; + let description = [{ + Lowered to `llvm::MDNode`. Operands are `#cir.md_string` or `#cir.md_node`. + + Example: + ```mlir + #cir.md_node<#cir.md_string<"agent">> + ``` + }]; + + let parameters = (ins OptionalArrayRefParameter<"mlir::Attribute">:$operands); + let assemblyFormat = [{ `<` (`>`) : ($operands^ `>`)? }]; + let genVerifyDecl = 1; + let canHaveIllegalCXXABIType = 0; +} + +def CIR_AnyMDAttr : AnyAttrOf<[CIR_MDStringAttr, CIR_MDNodeAttr]>; + //===----------------------------------------------------------------------===// // CIR_BlockAddrInfoAttr //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 985cc3712c581..f405e728aab31 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -4511,6 +4511,29 @@ def CIR_LocalInitOp : CIR_Op<"local_init", [ let hasLLVMLowering = false; } +//===----------------------------------------------------------------------===// +// MetadataAsValueOp +//===----------------------------------------------------------------------===// + +def CIR_MetadataAsValueOp : CIR_Op<"metadata_as_value", [Pure]> { + let summary = "Wrap a metadata attribute as an SSA value"; + let description = [{ + Provides `metadata` arguments to `cir.call_llvm_intrinsic`. Lowers to + `llvm.mlir.metadata_as_value`. + + Example: + + ```mlir + %0 = cir.metadata_as_value #cir.md_node<#cir.md_string<"agent">> + ``` + }]; + + let arguments = (ins CIR_AnyMDAttr:$metadata); + let results = (outs CIR_MetadataType:$res); + + let assemblyFormat = "$metadata attr-dict"; +} + //===----------------------------------------------------------------------===// // LLVMIntrinsicCallOp //===----------------------------------------------------------------------===// @@ -4520,14 +4543,15 @@ def CIR_LLVMIntrinsicCallOp : CIR_Op<"call_llvm_intrinsic"> { let description = [{ `cir.call_llvm_intrinsic` operation represents a call-like expression which has return type and arguments that maps directly to a llvm intrinsic. - It only records intrinsic `intrinsic_name`. + It only records intrinsic `intrinsic_name`. Arguments of LLVM `metadata` + type are passed as `!cir.metadata` values. }]; let results = (outs Optional<CIR_AnyType>:$result); let arguments = (ins StrAttr:$intrinsic_name, OptionalAttr<CIR_FastMathFlagsAttr>:$fastmath_flags, - Variadic<CIR_AnyType>:$arg_ops); + Variadic<AnyTypeOf<[CIR_AnyType, CIR_MetadataType]>>:$arg_ops); let skipDefaultBuilders = 1; diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index fd092c3ebb705..1f5c1613106f2 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td @@ -1224,6 +1224,18 @@ def CIR_CatchTokenType : CIR_Type<"CatchToken", "catch_token"> { }]; } +//===----------------------------------------------------------------------===// +// MetadataType +//===----------------------------------------------------------------------===// + +def CIR_MetadataType : CIR_Type<"Metadata", "metadata"> { + let summary = "CIR metadata type"; + let description = [{ + Mirrors the LLVM IR `metadata` type. Only valid as an operand of + `cir.call_llvm_intrinsic`. + }]; +} + //===----------------------------------------------------------------------===// // Global type constraints //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 89bfb4862c5c6..4a82a1c97cd75 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -141,6 +141,25 @@ emitAMDGCNImageOverloadedReturnType(CIRGenFunction &cgf, const CallExpr *e, return callOp.getResult(); } +static mlir::Value emitAMDGPUPredicate(CIRGenFunction &cgf, const CallExpr *e, + const llvm::Twine &name) { + CIRGenBuilderTy &builder = cgf.getBuilder(); + mlir::Location loc = cgf.getLoc(e->getExprLoc()); + mlir::MLIRContext *ctx = builder.getContext(); + + mlir::Value specId = + builder.getConstantInt(loc, builder.getUInt32Ty(), UINT32_MAX); + mlir::Value defaultValue = builder.getFalse(loc); + auto predicate = cir::MDNodeAttr::get( + ctx, cir::MDStringAttr::get(ctx, builder.getStringAttr(name))); + mlir::Value md = cir::MetadataAsValueOp::create(builder, loc, predicate); + mlir::Value result = builder.emitIntrinsicCallOp( + loc, "spv.named.boolean.spec.constant", builder.getBoolTy(), + mlir::ValueRange{specId, defaultValue, md}); + // Sema retypes the call to the logical operation type, which is int in C. + return builder.createBoolToInt(result, cgf.convertType(e->getType())); +} + std::optional<mlir::Value> CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, const CallExpr *expr) { @@ -602,6 +621,23 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, getContext().BuiltinInfo.getName(builtinId)); return mlir::Value{}; } + case AMDGPU::BI__builtin_amdgcn_processor_is: { + assert(cgm.getTriple().isSPIRV() && + "__builtin_amdgcn_processor_is should never reach CodeGen for " + "concrete targets!"); + StringRef proc = cast<clang::StringLiteral>(expr->getArg(0))->getString(); + return emitAMDGPUPredicate(*this, expr, "is." + proc); + } + case AMDGPU::BI__builtin_amdgcn_is_invocable: { + assert(cgm.getTriple().isSPIRV() && + "__builtin_amdgcn_is_invocable should never reach CodeGen for " + "concrete targets!"); + auto *fd = cast<FunctionDecl>( + cast<DeclRefExpr>(expr->getArg(0))->getReferencedDeclOfCallee()); + StringRef rf = + getContext().BuiltinInfo.getRequiredFeatures(fd->getBuiltinID()); + return emitAMDGPUPredicate(*this, expr, "has." + rf); + } case AMDGPU::BI__builtin_amdgcn_read_exec: case AMDGPU::BI__builtin_amdgcn_read_exec_lo: case AMDGPU::BI__builtin_amdgcn_read_exec_hi: { diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index 2649ed1449739..19ad998f079e3 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp @@ -1023,6 +1023,19 @@ LogicalResult EhFilterAttr::verify(function_ref<InFlightDiagnostic()> emitError, return success(); } +//===----------------------------------------------------------------------===// +// MDNodeAttr definitions +//===----------------------------------------------------------------------===// + +LogicalResult MDNodeAttr::verify(function_ref<InFlightDiagnostic()> emitError, + ArrayRef<mlir::Attribute> operands) { + for (mlir::Attribute operand : operands) + if (!mlir::isa<MDStringAttr, MDNodeAttr>(operand)) + return emitError() << "metadata node operands must be #cir.md_string or " + "#cir.md_node"; + return success(); +} + //===----------------------------------------------------------------------===// // RecordLayout lookup //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 6426841e6237f..17b90b4fbe1de 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -667,6 +667,25 @@ mlir::LogicalResult CIRToLLVMLLVMIntrinsicCallOpLowering::matchAndRewrite( return mlir::success(); } +static mlir::Attribute convertMetadataAttr(mlir::Attribute attr) { + mlir::MLIRContext *ctx = attr.getContext(); + if (auto str = mlir::dyn_cast<cir::MDStringAttr>(attr)) + return mlir::LLVM::MDStringAttr::get(ctx, str.getValue()); + auto node = mlir::cast<cir::MDNodeAttr>(attr); + SmallVector<mlir::Attribute> operands; + for (mlir::Attribute operand : node.getOperands()) + operands.push_back(convertMetadataAttr(operand)); + return mlir::LLVM::MDNodeAttr::get(ctx, operands); +} + +mlir::LogicalResult CIRToLLVMMetadataAsValueOpLowering::matchAndRewrite( + cir::MetadataAsValueOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const { + rewriter.replaceOpWithNewOp<mlir::LLVM::MetadataAsValueOp>( + op, convertMetadataAttr(op.getMetadata())); + return mlir::success(); +} + /// BoolAttr visitor. mlir::Value CIRAttrToValue::visitCirAttr(cir::BoolAttr boolAttr) { mlir::Location loc = parentOp->getLoc(); @@ -4157,6 +4176,9 @@ static void prepareTypeConverter(mlir::LLVMTypeConverter &converter, converter.addConversion([&](cir::VoidType type) -> mlir::Type { return mlir::LLVM::LLVMVoidType::get(type.getContext()); }); + converter.addConversion([&](cir::MetadataType type) -> mlir::Type { + return mlir::LLVM::LLVMMetadataType::get(type.getContext()); + }); } static void buildCtorDtorList( diff --git a/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c b/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c new file mode 100644 index 0000000000000..233f2c79bf4f2 --- /dev/null +++ b/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c @@ -0,0 +1,45 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR %s --input-file=%t.cir +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fclangir -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM %s --input-file=%t-cir.ll +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG %s --input-file=%t.ll + +void test_processor_is(void) { + if (__builtin_amdgcn_processor_is("gfx900")) + __builtin_trap(); +} + +// CIR-LABEL: cir.func {{.*}} @test_processor_is +// CIR: %[[ID:.*]] = cir.const #cir.int<4294967295> : !u32i +// CIR: %[[DEF:.*]] = cir.const #false +// CIR: %[[MD:.*]] = cir.metadata_as_value #cir.md_node<#cir.md_string<"is.gfx900">> +// CIR: cir.call_llvm_intrinsic "spv.named.boolean.spec.constant" %[[ID]], %[[DEF]], %[[MD]] : (!u32i, !cir.bool, !cir.metadata) -> !cir.bool + +// LLVM-LABEL: define {{.*}} @test_processor_is +// LLVM: call{{.*}} i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata ![[IS_GFX900:[0-9]+]]) + +// OGCG-LABEL: define {{.*}} @test_processor_is +// OGCG: call{{.*}} i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata ![[IS_GFX900:[0-9]+]]) + +void test_is_invocable(void) { + if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_permlanex16)) + __builtin_trap(); +} + +// CIR-LABEL: cir.func {{.*}} @test_is_invocable +// CIR: %[[MD:.*]] = cir.metadata_as_value #cir.md_node<#cir.md_string<"has.gfx10-insts">> +// CIR: cir.call_llvm_intrinsic "spv.named.boolean.spec.constant" %{{.*}}, %{{.*}}, %[[MD]] : (!u32i, !cir.bool, !cir.metadata) -> !cir.bool + +// LLVM-LABEL: define {{.*}} @test_is_invocable +// LLVM: call{{.*}} i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata ![[HAS_GFX10:[0-9]+]]) + +// OGCG-LABEL: define {{.*}} @test_is_invocable +// OGCG: call{{.*}} i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata ![[HAS_GFX10:[0-9]+]]) + +// LLVM-DAG: ![[IS_GFX900]] = !{!"is.gfx900"} +// LLVM-DAG: ![[HAS_GFX10]] = !{!"has.gfx10-insts"} + +// OGCG-DAG: ![[IS_GFX900]] = !{!"is.gfx900"} +// OGCG-DAG: ![[HAS_GFX10]] = !{!"has.gfx10-insts"} diff --git a/clang/test/CIR/IR/invalid-metadata.cir b/clang/test/CIR/IR/invalid-metadata.cir new file mode 100644 index 0000000000000..288aba61154d3 --- /dev/null +++ b/clang/test/CIR/IR/invalid-metadata.cir @@ -0,0 +1,21 @@ +// RUN: cir-opt %s -verify-diagnostics -split-input-file + +!s32i = !cir.int<s, 32> + +module { + cir.func @md_node_bad_operand() { + // expected-error@+1 {{metadata node operands must be #cir.md_string or #cir.md_node}} + %0 = cir.metadata_as_value #cir.md_node<#cir.int<0> : !s32i> + cir.return + } +} + +// ----- + +module { + cir.func @metadata_as_value_bad_attr() { + // expected-error@+1 {{attribute 'metadata' failed to satisfy constraint}} + %0 = cir.metadata_as_value #cir.bool<true> : !cir.bool + cir.return + } +} diff --git a/clang/test/CIR/Lowering/call-llvm-intrinsic.cir b/clang/test/CIR/Lowering/call-llvm-intrinsic.cir index 83893af858bd4..e4e5a3d3c8b16 100644 --- a/clang/test/CIR/Lowering/call-llvm-intrinsic.cir +++ b/clang/test/CIR/Lowering/call-llvm-intrinsic.cir @@ -45,3 +45,25 @@ module { cir.return %0 : !f32 } } + +module { + // Metadata operands are passed as !cir.metadata values. + // CHECK-LABEL: llvm.func @metadata_operand + // CHECK: %[[MD:.*]] = llvm.mlir.metadata_as_value #llvm.md_node<#llvm.md_string<"is.gfx900">> + // CHECK: llvm.call_intrinsic "llvm.spv.named.boolean.spec.constant"(%{{.*}}, %{{.*}}, %[[MD]]) : (i32, i1, !llvm.metadata) -> i1 + cir.func @metadata_operand(%arg0: !s32i, %arg1: !cir.bool) -> !cir.bool { + %0 = cir.metadata_as_value #cir.md_node<#cir.md_string<"is.gfx900">> + %1 = cir.call_llvm_intrinsic "spv.named.boolean.spec.constant" %arg0, %arg1, %0 : (!s32i, !cir.bool, !cir.metadata) -> !cir.bool + cir.return %1 : !cir.bool + } + + // CHECK-LABEL: llvm.func @metadata_nested + // CHECK: llvm.mlir.metadata_as_value #llvm.md_node<#llvm.md_string<"a">, #llvm.md_node<>> + // CHECK: llvm.mlir.metadata_as_value #llvm.md_string<"b"> + cir.func @metadata_nested() { + %0 = cir.metadata_as_value #cir.md_node<#cir.md_string<"a">, #cir.md_node<>> + %1 = cir.metadata_as_value #cir.md_string<"b"> + cir.call_llvm_intrinsic "donothing" : () -> () + cir.return + } +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
