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 <[email protected]> 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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
