[clang] [CIR][AMDGPU][SPIR-V] Support metadata intrinsic args for feature predicates (PR #229330)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 15 04:12:38 PDT 2026
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 <arseniy.obolenskiy at amd.com>
Date: Tue, 6 Oct 2026 09:41:32 +0200
Subject: [PATCH 1/3] [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 f6aaac5ab4cab1..e7e1f05f86dacf 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 985cc3712c581d..f405e728aab318 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 fd092c3ebb7056..1f5c1613106f23 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 89bfb4862c5c68..4a82a1c97cd757 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 2649ed14497393..19ad998f079e34 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 6426841e6237f3..17b90b4fbe1ded 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 00000000000000..233f2c79bf4f2d
--- /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 00000000000000..288aba61154d3f
--- /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 at +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 at +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 83893af858bd49..e4e5a3d3c8b164 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 <arseniy.obolenskiy at amd.com>
Date: Tue, 13 Oct 2026 05:56:53 +0200
Subject: [PATCH 2/3] 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 b6c37ddf1d97f6..eb4cf65b919c49 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 ac5b858c19576e..eef7edc594fb93 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 1f5c1613106f23..fd092c3ebb7056 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 bc2300aa1d067a..d58094052060e1 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 19ad998f079e34..2649ed14497393 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 299ecf2811cd03..b5010d5ece3111 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 233f2c79bf4f2d..d56789ced762c4 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 288aba61154d3f..00000000000000
--- 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 at +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 at +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 e4e5a3d3c8b164..83893af858bd49 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
- }
-}
>From 4349f95d014041cec9e58e4a0f8296ee79d727e9 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Thu, 15 Oct 2026 13:12:23 +0200
Subject: [PATCH 3/3] rename
---
clang/include/clang/CIR/Dialect/IR/CIROps.td | 8 ++++----
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 2 +-
clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 4 ++--
.../CodeGenBuiltins/builtin-amdgcn-feature-predicates.c | 4 ++--
4 files changed, 9 insertions(+), 9 deletions(-)
diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td
index eef7edc594fb93..536d570e3a9f0e 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIROps.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td
@@ -5623,11 +5623,11 @@ def CIR_AddrOfReturnAddrOp : CIR_Op<"address_of_return_address"> {
}
//===----------------------------------------------------------------------===//
-// AMDGCNFeaturePredicateOp
+// AMDGPUFeaturePredicateOp
//===----------------------------------------------------------------------===//
-def CIR_AMDGCNFeaturePredicateOp : CIR_Op<"amdgcn_feature_predicate"> {
- let summary = "AMDGCN feature predicate resolved at JIT time";
+def CIR_AMDGPUFeaturePredicateOp : CIR_Op<"amdgpu_feature_predicate"> {
+ let summary = "AMDGPU feature predicate resolved at JIT time";
let description = [{
Represents `__builtin_amdgcn_processor_is` and
@@ -5640,7 +5640,7 @@ def CIR_AMDGCNFeaturePredicateOp : CIR_Op<"amdgcn_feature_predicate"> {
Example:
```mlir
- %0 = cir.amdgcn_feature_predicate "is.gfx900"
+ %0 = cir.amdgpu_feature_predicate "is.gfx900"
```
}];
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index d58094052060e1..c18b6f03c0687d 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -145,7 +145,7 @@ 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::Value result = cir::AMDGCNFeaturePredicateOp::create(
+ mlir::Value result = cir::AMDGPUFeaturePredicateOp::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/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index b5010d5ece3111..25d759be04d738 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -2452,8 +2452,8 @@ mlir::LogicalResult CIRToLLVMAddrOfReturnAddrOpLowering::matchAndRewrite(
return mlir::success();
}
-mlir::LogicalResult CIRToLLVMAMDGCNFeaturePredicateOpLowering::matchAndRewrite(
- cir::AMDGCNFeaturePredicateOp op, OpAdaptor adaptor,
+mlir::LogicalResult CIRToLLVMAMDGPUFeaturePredicateOpLowering::matchAndRewrite(
+ cir::AMDGPUFeaturePredicateOp op, OpAdaptor adaptor,
mlir::ConversionPatternRewriter &rewriter) const {
mlir::Location loc = op.getLoc();
mlir::MLIRContext *ctx = rewriter.getContext();
diff --git a/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c b/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c
index d56789ced762c4..4088ee1e1cf22b 100644
--- a/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c
+++ b/clang/test/CIR/CodeGenBuiltins/builtin-amdgcn-feature-predicates.c
@@ -12,7 +12,7 @@ void test_processor_is(void) {
}
// CIR-LABEL: cir.func {{.*}} @test_processor_is
-// CIR: cir.amdgcn_feature_predicate "is.gfx900"
+// CIR: cir.amdgpu_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]+]])
@@ -26,7 +26,7 @@ void test_is_invocable(void) {
}
// CIR-LABEL: cir.func {{.*}} @test_is_invocable
-// CIR: cir.amdgcn_feature_predicate "has.gfx10-insts"
+// CIR: cir.amdgpu_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]+]])
More information about the cfe-commits
mailing list