[clang] 4812a67 - [CIR][AMDGPU] Add support for AMDGCN frexp_mant builtins (#198121)

[via cfe-commits]

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);
+}