[clang] [CIR][AMDGPU] Use intrinsic type for dpp signature detection (PR #226921)
Steffen Larsen via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 2 00:11:11 PDT 2026
https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/226921
>From d6ca216653256320caccd3e53b7375bcd2e16f99 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Mon, 28 Sep 2026 02:46:39 -0500
Subject: [PATCH] [CIR][AMDGPU] Use intrinsic type for dpp signature detection
This commit modifies the CIR code generation for AMDGPU to use the
intrinsic type for dpp builtin signature detection. This change ensures
that the correct intrinsic type is used when generating code for dpp
operations, instead of relying on hard-coded types.
This should also help simplify the implementation of other future
builtins, such as NVPTX wmma/mma load/store/mma.
Note that this is not an RFC as it changes the signedness in parts of
the builtin signatures. However, that change fits the signature
definitions and the signedness does not carry much meaning for these
builtins either way.
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 24 ++++++++++++-------
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 21 +++++-----------
clang/lib/CIR/CodeGen/CIRGenFunction.h | 6 +++++
.../CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip | 12 +++++-----
4 files changed, 33 insertions(+), 30 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
index 245708691b7d9..ebed52080cf1a 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
@@ -1132,13 +1132,19 @@ static RValue tryEmitFPMathIntrinsic(CIRGenFunction &cgf, const CallExpr *e,
static mlir::Type
decodeFixedType(CIRGenFunction &cgf,
ArrayRef<llvm::Intrinsic::IITDescriptor> &infos,
- mlir::MLIRContext *context) {
+ ArrayRef<mlir::Type> overloadTys, mlir::MLIRContext *context) {
using namespace llvm::Intrinsic;
IITDescriptor descriptor = infos.front();
infos = infos.slice(1);
switch (descriptor.Kind) {
+ case IITDescriptor::Overloaded:
+ case IITDescriptor::Match:
+ if (descriptor.getOverloadIndex() < overloadTys.size())
+ return overloadTys[descriptor.getOverloadIndex()];
+ cgf.cgm.errorNYI("Overloaded intrinsic type without overload types");
+ return cir::VoidType::get(context);
case IITDescriptor::Void:
return cir::VoidType::get(context);
case IITDescriptor::Half:
@@ -1157,7 +1163,7 @@ decodeFixedType(CIRGenFunction &cgf,
return cir::IntType::get(context, descriptor.IntegerWidth,
/*isSigned=*/true);
case IITDescriptor::Vector: {
- mlir::Type elementType = decodeFixedType(cgf, infos, context);
+ mlir::Type elementType = decodeFixedType(cgf, infos, overloadTys, context);
unsigned numElements = descriptor.VectorWidth.getFixedValue();
return cir::VectorType::get(elementType, numElements);
}
@@ -1227,19 +1233,20 @@ static mlir::Value getCorrectedPtr(mlir::Value argValue, mlir::Type expectedTy,
return builder.createBitcast(argValue, expectedTy);
}
-static cir::FuncType getIntrinsicType(CIRGenFunction &cgf,
- mlir::MLIRContext *context,
- llvm::Intrinsic::ID id) {
+cir::FuncType
+CIRGenFunction::getIntrinsicType(llvm::Intrinsic::ID id,
+ ArrayRef<mlir::Type> overloadTys) {
using namespace llvm::Intrinsic;
+ mlir::MLIRContext *context = &getMLIRContext();
SmallVector<IITDescriptor, 8> table;
auto [tableRef, _, isVarArg] = getIntrinsicInfoTableEntries(id, table);
- mlir::Type resultTy = decodeFixedType(cgf, tableRef, context);
+ mlir::Type resultTy = decodeFixedType(*this, tableRef, overloadTys, context);
SmallVector<mlir::Type, 8> argTypes;
while (!tableRef.empty())
- argTypes.push_back(decodeFixedType(cgf, tableRef, context));
+ argTypes.push_back(decodeFixedType(*this, tableRef, overloadTys, context));
// CIR convention: no explicit void return type
if (isa<cir::VoidType>(resultTy))
@@ -3188,8 +3195,7 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID,
assert(name.starts_with("llvm.") && "expected llvm. prefix");
name = name.drop_front(/*strlen("llvm.")=*/5);
- cir::FuncType intrinsicType =
- getIntrinsicType(*this, &getMLIRContext(), intrinsicID);
+ cir::FuncType intrinsicType = getIntrinsicType(intrinsicID);
SmallVector<mlir::Value> args;
const FunctionDecl *fd = e->getDirectCallee();
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index b2aac0d376fa1..13fa6fc613783 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -14,6 +14,7 @@
#include "mlir/IR/Value.h"
#include "clang/Basic/TargetBuiltins.h"
+#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
#include "llvm/Support/ErrorHandling.h"
@@ -223,15 +224,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
bool isUpdateDpp = builtinId == AMDGPU::BI__builtin_amdgcn_update_dpp;
llvm::StringRef intrinsicName =
isMovDpp8 ? "amdgcn.mov.dpp8" : "amdgcn.update.dpp";
-
- // Fixed parameter types of the target LLVM intrinsics, following the
- // "old"/"data" operands which share the overloaded integer type.
- llvm::SmallVector<mlir::Type, 4> fixedTailTypes;
- mlir::Type ui32 = builder.getUInt32Ty();
- if (isMovDpp8)
- fixedTailTypes = {ui32};
- else
- fixedTailTypes = {ui32, ui32, ui32, builder.getUIntNTy(1)};
+ cir::FuncType intrinsicTy =
+ getIntrinsicType(isMovDpp8 ? llvm::Intrinsic::amdgcn_mov_dpp8
+ : llvm::Intrinsic::amdgcn_update_dpp,
+ {intTy});
auto coerceTo = [&](mlir::Value from, mlir::Type to) -> mlir::Value {
if (from.getType() == to)
@@ -252,7 +248,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
// Number of builtin-level leading args that need zero-extend promotion when
// the data type is narrower than 32 bits.
unsigned numPromotedArgs = isUpdateDpp ? 2u : 1u;
- unsigned numIntTyFinalPos = isMovDpp8 ? 1u : 2u;
for (unsigned i = 0; i != expr->getNumArgs(); ++i) {
mlir::Value v =
emitScalarOrConstFoldImmArg(iceArguments, i, expr->getArg(i));
@@ -262,11 +257,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
v = builder.createBitcast(v, sameWidthUTy);
v = builder.createIntCast(v, intTy);
}
- unsigned finalIdx = i + unsigned(isMovDpp);
- mlir::Type finalTy = finalIdx < numIntTyFinalPos
- ? intTy
- : fixedTailTypes[finalIdx - numIntTyFinalPos];
- args.push_back(coerceTo(v, finalTy));
+ args.push_back(coerceTo(v, intrinsicTy.getInput(i + unsigned(isMovDpp))));
}
mlir::Value result = builder.emitIntrinsicCallOp(loc, intrinsicName, intTy,
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index b232abf5d9299..26f043e472299 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -38,6 +38,7 @@
#include "clang/CIR/TypeEvaluationKind.h"
#include "llvm/ADT/ScopedHashTable.h"
#include "llvm/IR/Instructions.h"
+#include "llvm/IR/Intrinsics.h"
namespace {
class ScalarExprEmitter;
@@ -1917,6 +1918,11 @@ class CIRGenFunction : public CIRGenTypeCache {
int64_t getAccessedFieldNo(unsigned idx, mlir::ArrayAttr elts);
+ /// Return the CIR signature of the LLVM intrinsic \p id, resolving its
+ /// overloaded types to \p overloadTys. Integer types are returned signed.
+ cir::FuncType getIntrinsicType(llvm::Intrinsic::ID id,
+ llvm::ArrayRef<mlir::Type> overloadTys = {});
+
/// Emit a simple LLVM intrinsic that takes N scalar arguments. The intrinsic
/// name is used verbatim; any overload mangling (e.g. `.f32`, `.p1`) must be
/// baked into \p intrinName by the caller. The result type defaults to the
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
index e15c8fbc59944..8b7f4dbf9c0bd 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
@@ -59,7 +59,7 @@ __device__ void test_permlanex16(unsigned int* out, unsigned int a, unsigned int
// CIR-LABEL: @_Z16test_mov_dpp_intPii
// CIR: %[[OLD:.*]] = cir.const #cir.poison : !u32i
-// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" %[[OLD]], {{.*}} : (!u32i, !u32i, !u32i, !u32i, !u32i, !cir.int<u, 1>) -> !u32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" %[[OLD]], {{.*}} : (!u32i, !u32i, !s32i, !s32i, !s32i, !cir.int<s, 1>) -> !u32i
// LLVM: define{{.*}} void @_Z16test_mov_dpp_intPii
// LLVM: call i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %{{.*}}, i32 257, i32 15, i32 15, i1 false)
__device__ void test_mov_dpp_int(int* out, int src) {
@@ -68,7 +68,7 @@ __device__ void test_mov_dpp_int(int* out, int src) {
// CIR-LABEL: @_Z18test_mov_dpp_shortsPs
// CIR: cir.cast bitcast {{.*}} : !s16i -> !u16i
-// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !u32i, !cir.int<u, 1>) -> !u32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" {{.*}} : (!u32i, !u32i, !s32i, !s32i, !s32i, !cir.int<s, 1>) -> !u32i
// LLVM: define{{.*}} void @_Z18test_mov_dpp_shortsPs
// LLVM: call i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %{{.*}}, i32 257, i32 15, i32 15, i1 false)
__device__ void test_mov_dpp_short(short x, short *p) {
@@ -77,7 +77,7 @@ __device__ void test_mov_dpp_short(short x, short *p) {
// CIR-LABEL: @_Z18test_mov_dpp_floatfPf
// CIR: cir.cast bitcast {{.*}} : !cir.float -> !u32i
-// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !u32i, !cir.int<u, 1>) -> !u32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" {{.*}} : (!u32i, !u32i, !s32i, !s32i, !s32i, !cir.int<s, 1>) -> !u32i
// LLVM: define{{.*}} void @_Z18test_mov_dpp_floatfPf
// LLVM: call i32 @llvm.amdgcn.update.dpp.i32(i32 poison, i32 %{{.*}}, i32 257, i32 15, i32 15, i1 false)
__device__ void test_mov_dpp_float(float x, float *p) {
@@ -85,7 +85,7 @@ __device__ void test_mov_dpp_float(float x, float *p) {
}
// CIR-LABEL: @_Z19test_update_dpp_intPiii
-// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !u32i, !cir.int<u, 1>) -> !u32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" {{.*}} : (!u32i, !u32i, !s32i, !s32i, !s32i, !cir.int<s, 1>) -> !u32i
// LLVM: define{{.*}} void @_Z19test_update_dpp_intPiii
// LLVM: call i32 @llvm.amdgcn.update.dpp.i32(i32 %{{.*}}, i32 %{{.*}}, i32 0, i32 0, i32 0, i1 false)
__device__ void test_update_dpp_int(int* out, int arg1, int arg2) {
@@ -93,7 +93,7 @@ __device__ void test_update_dpp_int(int* out, int arg1, int arg2) {
}
// CIR-LABEL: @_Z18test_mov_dpp8_uintPjj
-// CIR: cir.call_llvm_intrinsic "amdgcn.mov.dpp8" {{.*}} : (!u32i, !u32i) -> !u32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.mov.dpp8" {{.*}} : (!u32i, !s32i) -> !u32i
// LLVM: define{{.*}} void @_Z18test_mov_dpp8_uintPjj
// LLVM: call i32 @llvm.amdgcn.mov.dpp8.i32(i32 %{{.*}}, i32 1)
__device__ void test_mov_dpp8_uint(unsigned int* out, unsigned int a) {
@@ -102,7 +102,7 @@ __device__ void test_mov_dpp8_uint(unsigned int* out, unsigned int a) {
// CIR-LABEL: @_Z19test_mov_dpp8_shortsPs
// CIR: cir.cast bitcast {{.*}} : !s16i -> !u16i
-// CIR: cir.call_llvm_intrinsic "amdgcn.mov.dpp8" {{.*}} : (!u32i, !u32i) -> !u32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.mov.dpp8" {{.*}} : (!u32i, !s32i) -> !u32i
// LLVM: define{{.*}} void @_Z19test_mov_dpp8_shortsPs
// LLVM: call i32 @llvm.amdgcn.mov.dpp8.i32(i32 %{{.*}}, i32 1)
__device__ void test_mov_dpp8_short(short x, short *p) {
More information about the cfe-commits
mailing list