[clang] [CIR][AMDGPU] Add support for AMDGCN frexp_mant builtins (PR #198121)

[Ayokunle Amodu via cfe-commits]

https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/198121

>From 99c849226928c23bcb4bd01963faa3bef6ccadb5 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sat, 16 May 2026 16:15:16 -0400
Subject: [PATCH 1/3] add amdgcn frexp mant buitins

---
 clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp    |  5 +----
 .../CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip    |  8 ++++++++
 clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip    | 16 ++++++++++++++++
 3 files changed, 25 insertions(+), 4 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0a7ba0c194400..f74379db0be54 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -330,10 +330,7 @@ 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");
   }
   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 c21359ff7a88b..bb0c95236ab2d 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -63,3 +63,11 @@
 __device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) {
   *out = __builtin_amdgcn_div_fixuph(a, b, c);
 }
+
+// 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);
+}
\ No newline at end of file
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index ca708bca8587b..b0eecfb7cc755 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -111,3 +111,19 @@ __device__ void test_readfirstlane(int* out, int a) {
 __device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out) {
   *out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr();
 }
+
+// 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);
+}

>From bd506d13d5a40e61faba3428891d10dca3f3d82b Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 12:15:41 -0400
Subject: [PATCH 2/3] add nl to test

---
 clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
index bb0c95236ab2d..987a8c27aa4d3 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -70,4 +70,4 @@ __device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float
 // 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);
-}
\ No newline at end of file
+}

>From 5e05252735d7327d710005fdd8c0b432130d3176 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 12:27:13 -0400
Subject: [PATCH 3/3] return correct value for builtin

---
 clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index f74379db0be54..3b860b11924f4 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -330,7 +330,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: {
-    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.frexp.mant");
+    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.frexp.mant")
+        .getValue();
   }
   case AMDGPU::BI__builtin_amdgcn_frexp_exp:
   case AMDGPU::BI__builtin_amdgcn_frexp_expf: