Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

[CIR][AMDGPU] Add support for AMDGCN ldexp builtins (#198087)

Browse source 
Adds support for the following AMDGCN ldexp builtins:

- __builtin_amdgcn_ldexp (double)
- __builtin_amdgcn_ldexpf (float)
- __builtin_amdgcn_ldexph (half)

These are lowered to the `llvm.ldexp` intrinsic.
This commit is contained in:
Ayokunle Amodu 2026-07-31 19:18:49 -04:00 committed by GitHub
parent 362d72f638
commit b0511699c3
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
3 changed files with 40 additions and 4 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

18
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
View file

@ -296,10 +296,20 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_ldexp:
case AMDGPU::BI__builtin_amdgcn_ldexpf:
case AMDGPU::BI__builtin_amdgcn_ldexph: {
cgm.errorNYI(expr->getSourceRange(),
std::string("unimplemented AMDGPU builtin call: ") +
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
mlir::Value src0 = emitScalarExpr(expr->getArg(0));
mlir::Value src1 = emitScalarExpr(expr->getArg(1));
// For ldexph, the raw instruction has different behavior for out-of-bounds
// exponent values (implicit truncation instead of saturate to
// short_min/short_max), so truncate the exponent to i16 first.
mlir::Value exp =
builtinId == AMDGPU::BI__builtin_amdgcn_ldexph
? cir::CastOp::create(builder, getLoc(expr->getExprLoc()),
builder.getSInt16Ty(),
cir::CastKind::integral, src1)
: src1;
return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()), "ldexp",
src0.getType(),
mlir::ValueRange{src0, exp});
}
case AMDGPU::BI__builtin_amdgcn_frexp_mant:
case AMDGPU::BI__builtin_amdgcn_frexp_mantf:

10
clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
View file

@ -111,3 +111,13 @@ __device__ void test_cos_f16(_Float16* out, _Float16 a) {
__device__ void test_frexp_mant_f16(_Float16* out, _Float16 a) {
*out = __builtin_amdgcn_frexp_manth(a);
}
// CIR-LABEL: @_Z14test_ldexp_f16PDF16_DF16_i
// CIR: [[TRUNC:%.*]] = cir.cast integral {{.*}} : !s32i -> !s16i
// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}}, [[TRUNC]] : (!cir.f16, !s16i) -> !cir.f16
// LLVM: define{{.*}} void @_Z14test_ldexp_f16PDF16_DF16_i
// LLVM: [[TRUNC:%.*]] = trunc i32 {{.*}} to i16
// LLVM: call{{.*}} half @llvm.ldexp.f16.i16(half %{{.*}}, i16 [[TRUNC]])
__device__ void test_ldexp_f16(_Float16* out, _Float16 a, int b) {
*out = __builtin_amdgcn_ldexph(a, b);
}

16
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
View file

@ -239,3 +239,19 @@ __device__ void test_trig_preop_f32(float* out, float a, int b) {
__device__ void test_trig_preop_f64(double* out, double a, int b) {
*out = __builtin_amdgcn_trig_preop(a, b);
}
// CIR-LABEL: @_Z14test_ldexp_f32Pffi
// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.float, !s32i) -> !cir.float
// LLVM: define{{.*}} void @_Z14test_ldexp_f32Pffi
// LLVM: call{{.*}} float @llvm.ldexp.f32.i32(float %{{.*}}, i32 %{{.*}})
__device__ void test_ldexp_f32(float* out, float a, int b) {
*out = __builtin_amdgcn_ldexpf(a, b);
}
// CIR-LABEL: @_Z14test_ldexp_f64Pddi
// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.double, !s32i) -> !cir.double
// LLVM: define{{.*}} void @_Z14test_ldexp_f64Pddi
// LLVM: call{{.*}} double @llvm.ldexp.f64.i32(double %{{.*}}, i32 %{{.*}})
__device__ void test_ldexp_f64(double* out, double a, int b) {
*out = __builtin_amdgcn_ldexp(a, b);
}