https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/229330
>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 1/2] [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 + } +} >From 88663239fff1618d24959694c8ad4ee659e4d062 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Tue, 13 Oct 2026 05:56:53 +0200 Subject: [PATCH 2/2] Address comments --- .../include/clang/CIR/Dialect/IR/CIRAttrs.td | 39 ------------- clang/include/clang/CIR/Dialect/IR/CIROps.td | 55 ++++++++++--------- .../include/clang/CIR/Dialect/IR/CIRTypes.td | 12 ---- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 13 +---- clang/lib/CIR/Dialect/IR/CIRAttrs.cpp | 13 ----- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 42 +++++++------- .../builtin-amdgcn-feature-predicates.c | 8 +-- clang/test/CIR/IR/invalid-metadata.cir | 21 ------- .../test/CIR/Lowering/call-llvm-intrinsic.cir | 22 -------- 9 files changed, 53 insertions(+), 172 deletions(-) delete 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 b6c37ddf1d97f..eb4cf65b919c4 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -1970,45 +1970,6 @@ 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 ac5b858c19576..eef7edc594fb9 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -4513,29 +4513,6 @@ 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"; -} - //===----------------------------------------------------------------------===// // RegisterExitDtorOp //===----------------------------------------------------------------------===// @@ -4590,15 +4567,14 @@ 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`. Arguments of LLVM `metadata` - type are passed as `!cir.metadata` values. + It only records intrinsic `intrinsic_name`. }]; let results = (outs Optional<CIR_AnyType>:$result); let arguments = (ins StrAttr:$intrinsic_name, OptionalAttr<CIR_FastMathFlagsAttr>:$fastmath_flags, - Variadic<AnyTypeOf<[CIR_AnyType, CIR_MetadataType]>>:$arg_ops); + Variadic<CIR_AnyType>:$arg_ops); let skipDefaultBuilders = 1; @@ -5646,6 +5622,33 @@ def CIR_AddrOfReturnAddrOp : CIR_Op<"address_of_return_address"> { let assemblyFormat = "attr-dict `:` qualified(type($result))"; } +//===----------------------------------------------------------------------===// +// AMDGCNFeaturePredicateOp +//===----------------------------------------------------------------------===// + +def CIR_AMDGCNFeaturePredicateOp : CIR_Op<"amdgcn_feature_predicate"> { + let summary = "AMDGCN feature predicate resolved at JIT time"; + + let description = [{ + Represents `__builtin_amdgcn_processor_is` and + `__builtin_amdgcn_is_invocable` when targeting AMDGCN-flavored SPIR-V. + The result is a specialization constant resolved once the concrete + processor is known. + + `predicate` is `is.<processor>` or `has.<features>`. + + Example: + + ```mlir + %0 = cir.amdgcn_feature_predicate "is.gfx900" + ``` + }]; + + let arguments = (ins StrAttr:$predicate); + let results = (outs CIR_BoolType:$result); + let assemblyFormat = "$predicate attr-dict"; +} + //===----------------------------------------------------------------------===// // StackSaveOp & StackRestoreOp //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td index 1f5c1613106f2..fd092c3ebb705 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRTypes.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRTypes.td @@ -1224,18 +1224,6 @@ 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 bc2300aa1d067..d58094052060e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -145,17 +145,8 @@ 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}); + mlir::Value result = cir::AMDGCNFeaturePredicateOp::create( + builder, loc, builder.getBoolTy(), builder.getStringAttr(name)); // Sema retypes the call to the logical operation type, which is int in C. return builder.createBoolToInt(result, cgf.convertType(e->getType())); } diff --git a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp index 19ad998f079e3..2649ed1449739 100644 --- a/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRAttrs.cpp @@ -1023,19 +1023,6 @@ 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 299ecf2811cd0..b5010d5ece311 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -667,25 +667,6 @@ 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(); @@ -2471,6 +2452,26 @@ mlir::LogicalResult CIRToLLVMAddrOfReturnAddrOpLowering::matchAndRewrite( return mlir::success(); } +mlir::LogicalResult CIRToLLVMAMDGCNFeaturePredicateOpLowering::matchAndRewrite( + cir::AMDGCNFeaturePredicateOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const { + mlir::Location loc = op.getLoc(); + mlir::MLIRContext *ctx = rewriter.getContext(); + mlir::Value specId = + mlir::LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(), -1); + mlir::Value defaultValue = + mlir::LLVM::ConstantOp::create(rewriter, loc, rewriter.getI1Type(), 0); + auto predicate = mlir::LLVM::MDNodeAttr::get( + ctx, mlir::LLVM::MDStringAttr::get(ctx, op.getPredicateAttr())); + mlir::Value md = + mlir::LLVM::MetadataAsValueOp::create(rewriter, loc, predicate); + replaceOpWithCallLLVMIntrinsicOp( + rewriter, op, "llvm.spv.named.boolean.spec.constant", + getTypeConverter()->convertType(op.getType()), + mlir::ValueRange{specId, defaultValue, md}); + return mlir::success(); +} + mlir::LogicalResult CIRToLLVMLoadOpLowering::matchAndRewrite( cir::LoadOp op, OpAdaptor adaptor, mlir::ConversionPatternRewriter &rewriter) const { @@ -4196,9 +4197,6 @@ 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( mlir::ModuleOp module, StringRef globalXtorName, StringRef llvmXtorName, diff --git a/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c b/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c index 233f2c79bf4f2..d56789ced762c 100644 --- a/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c +++ b/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c @@ -12,10 +12,7 @@ void test_processor_is(void) { } // 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 +// CIR: cir.amdgcn_feature_predicate "is.gfx900" // LLVM-LABEL: define {{.*}} @test_processor_is // LLVM: call{{.*}} i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata ![[IS_GFX900:[0-9]+]]) @@ -29,8 +26,7 @@ void test_is_invocable(void) { } // 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 +// CIR: cir.amdgcn_feature_predicate "has.gfx10-insts" // LLVM-LABEL: define {{.*}} @test_is_invocable // LLVM: call{{.*}} i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata ![[HAS_GFX10:[0-9]+]]) diff --git a/clang/test/CIR/IR/invalid-metadata.cir b/clang/test/CIR/IR/invalid-metadata.cir deleted file mode 100644 index 288aba61154d3..0000000000000 --- a/clang/test/CIR/IR/invalid-metadata.cir +++ /dev/null @@ -1,21 +0,0 @@ -// 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 e4e5a3d3c8b16..83893af858bd4 100644 --- a/clang/test/CIR/Lowering/call-llvm-intrinsic.cir +++ b/clang/test/CIR/Lowering/call-llvm-intrinsic.cir @@ -45,25 +45,3 @@ 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
