Author: Ayokunle Amodu Date: 2026-07-18T12:57:05-04:00 New Revision: 4812a67a45710531f9ac3bf2a84a1e2b278971a4
URL: https://github.com/llvm/llvm-project/commit/4812a67a45710531f9ac3bf2a84a1e2b278971a4 DIFF: https://github.com/llvm/llvm-project/commit/4812a67a45710531f9ac3bf2a84a1e2b278971a4.diff LOG: [CIR][AMDGPU] Add support for AMDGCN frexp_mant builtins (#198121) Adds codegen for the following AMDGCN frexp mantissa builtins: - __builtin_amdgcn_frexp_mant (double) - __builtin_amdgcn_frexp_mantf (float) - __builtin_amdgcn_frexp_manth (half) These are lowered to the corresponding `llvm.amdgcn.frexp.mant` intrinsic. Added: Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 43ca7cbc11de5..63bc9d8b6b144 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -307,10 +307,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_frexp_mant: case AMDGPU::BI__builtin_amdgcn_frexp_mantf: case AMDGPU::BI__builtin_amdgcn_frexp_manth: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.frexp.mant") + .getValue(); } case AMDGPU::BI__builtin_amdgcn_frexp_exp: case AMDGPU::BI__builtin_amdgcn_frexp_expf: diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip index ce5168996ab28..e7c4a2db9789f 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip @@ -103,3 +103,11 @@ __device__ void test_sin_f16(_Float16* out, _Float16 a) { __device__ void test_cos_f16(_Float16* out, _Float16 a) { *out = __builtin_amdgcn_cosh(a); } + +// CIR-LABEL: @_Z19test_frexp_mant_f16PDF16_DF16_ +// CIR: cir.call_llvm_intrinsic "amdgcn.frexp.mant" {{.*}} : (!cir.f16) -> !cir.f16 +// LLVM: define{{.*}} void @_Z19test_frexp_mant_f16PDF16_DF16_ +// LLVM: call{{.*}} half @llvm.amdgcn.frexp.mant.f16(half %{{.*}}) +__device__ void test_frexp_mant_f16(_Float16* out, _Float16 a) { + *out = __builtin_amdgcn_frexp_manth(a); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 184a6ec7a146c..5dfc95490209b 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -207,3 +207,19 @@ __device__ void test_exp2_f32(float* out, float a) { __device__ void test_log_f32(float* out, float a) { *out = __builtin_amdgcn_logf(a); } + +// CIR-LABEL: @_Z19test_frexp_mant_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.frexp.mant" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z19test_frexp_mant_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.frexp.mant.f32(float %{{.*}}) +__device__ void test_frexp_mant_f32(float* out, float a) { + *out = __builtin_amdgcn_frexp_mantf(a); +} + +// CIR-LABEL: @_Z19test_frexp_mant_f64Pdd +// CIR: cir.call_llvm_intrinsic "amdgcn.frexp.mant" {{.*}} : (!cir.double) -> !cir.double +// LLVM: define{{.*}} void @_Z19test_frexp_mant_f64Pdd +// LLVM: call{{.*}} double @llvm.amdgcn.frexp.mant.f64(double %{{.*}}) +__device__ void test_frexp_mant_f64(double* out, double a) { + *out = __builtin_amdgcn_frexp_mant(a); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
