[clang] 084a4e6 - [CIR][AMDGPU] Implement __builtin_amdgcn_*_dpp* builtins (#226469)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 28 01:14:28 PDT 2026
Author: Steffen Larsen
Date: 2026-09-28T10:14:22+02:00
New Revision: 084a4e6002a9b718adb31dd906b233cafd3a9814
URL: https://github.com/llvm/llvm-project/commit/084a4e6002a9b718adb31dd906b233cafd3a9814
DIFF: https://github.com/llvm/llvm-project/commit/084a4e6002a9b718adb31dd906b233cafd3a9814.diff
LOG: [CIR][AMDGPU] Implement __builtin_amdgcn_*_dpp* builtins (#226469)
This commit implements the `__builtin_amdgcn_update_dpp`,
`__builtin_amdgcn_mov_dpp`, and `__builtin_amdgcn_mov_dpp8` builtins in
CIR, closely matching the implementation in
CodeGenFunction::EmitAMDGPUBuiltinExpr from OGCG.
Assisted-by: Claude Sonnet 5
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
Added:
Modified:
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
Removed:
################################################################################
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 58ef4b1bd4b61..b2aac0d376fa1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -206,10 +206,74 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
case AMDGPU::BI__builtin_amdgcn_mov_dpp:
case AMDGPU::BI__builtin_amdgcn_update_dpp: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ mlir::Location loc = getLoc(expr->getExprLoc());
+ unsigned iceArguments = 0;
+ ASTContext::GetBuiltinTypeError error;
+ getContext().GetBuiltinType(builtinId, error, &iceArguments);
+ assert(error == ASTContext::GE_None && "Should not codegen an error");
+ assert(expr->getNumArgs() == 5 || expr->getNumArgs() == 6 ||
+ expr->getNumArgs() == 2);
+
+ mlir::Type dataTy = convertType(expr->getArg(0)->getType());
+ unsigned size = cgm.getDataLayout().getTypeSizeInBits(dataTy);
+ cir::IntType intTy = builder.getUIntNTy(std::max(size, 32u));
+
+ bool isMovDpp8 = builtinId == AMDGPU::BI__builtin_amdgcn_mov_dpp8;
+ bool isMovDpp = builtinId == AMDGPU::BI__builtin_amdgcn_mov_dpp;
+ 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)};
+
+ auto coerceTo = [&](mlir::Value from, mlir::Type to) -> mlir::Value {
+ if (from.getType() == to)
+ return from;
+ if (mlir::isa<cir::IntType>(from.getType()) &&
+ mlir::isa<cir::IntType>(to))
+ return builder.createIntCast(from, to);
+ return builder.createBitcast(from, to);
+ };
+
+ llvm::SmallVector<mlir::Value, 6> args;
+ // __builtin_amdgcn_mov_dpp has no "old" operand at the source level, but
+ // the real intrinsic it lowers to requires one, so we synthesize a poison
+ // value for it since it is never meaningfully read.
+ if (isMovDpp)
+ args.push_back(builder.getConstant(loc, cir::PoisonAttr::get(intTy)));
+
+ // 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));
+ if (i < numPromotedArgs && size < 32) {
+ mlir::Type sameWidthUTy = builder.getUIntNTy(size);
+ if (v.getType() != sameWidthUTy)
+ 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));
+ }
+
+ mlir::Value result = builder.emitIntrinsicCallOp(loc, intrinsicName, intTy,
+ mlir::ValueRange(args));
+ if (size < 32 && !mlir::isa<cir::IntType>(dataTy))
+ result = builder.createIntCast(result, builder.getUIntNTy(size));
+ return coerceTo(result, dataTy);
}
case AMDGPU::BI__builtin_amdgcn_permlane16:
case AMDGPU::BI__builtin_amdgcn_permlanex16: {
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
index d3ce74b00e093..e15c8fbc59944 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx10.hip
@@ -56,3 +56,55 @@ __device__ void test_permlane16(unsigned int* out, unsigned int a, unsigned int
__device__ void test_permlanex16(unsigned int* out, unsigned int a, unsigned int b, unsigned int c, unsigned int d) {
*out = __builtin_amdgcn_permlanex16(a, b, c, d, 0, 0);
}
+
+// 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
+// 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) {
+ *out = __builtin_amdgcn_mov_dpp(src, 0x101, 0xf, 0xf, 0);
+}
+
+// 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
+// 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) {
+ *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
+}
+
+// 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
+// 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) {
+ *p = __builtin_amdgcn_mov_dpp(x, 0x101, 0xf, 0xf, 0);
+}
+
+// CIR-LABEL: @_Z19test_update_dpp_intPiii
+// CIR: cir.call_llvm_intrinsic "amdgcn.update.dpp" {{.*}} : (!u32i, !u32i, !u32i, !u32i, !u32i, !cir.int<u, 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) {
+ *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
+}
+
+// CIR-LABEL: @_Z18test_mov_dpp8_uintPjj
+// CIR: cir.call_llvm_intrinsic "amdgcn.mov.dpp8" {{.*}} : (!u32i, !u32i) -> !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) {
+ *out = __builtin_amdgcn_mov_dpp8(a, 1);
+}
+
+// CIR-LABEL: @_Z19test_mov_dpp8_shortsPs
+// CIR: cir.cast bitcast {{.*}} : !s16i -> !u16i
+// CIR: cir.call_llvm_intrinsic "amdgcn.mov.dpp8" {{.*}} : (!u32i, !u32i) -> !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) {
+ *p = __builtin_amdgcn_mov_dpp8(x, 1);
+}
More information about the cfe-commits
mailing list