[clang] 00e7c35 - [CIR][AMDGPU] Add support for AMDGCN tanh builtins (#197852)

[via cfe-commits]

Author: Ayokunle Amodu
Date: 2026-07-31T14:57:40-04:00
New Revision: 00e7c3564c2ae34d4ac2c6aaa838ac917472ed23

URL: https://github.com/llvm/llvm-project/commit/00e7c3564c2ae34d4ac2c6aaa838ac917472ed23
DIFF: https://github.com/llvm/llvm-project/commit/00e7c3564c2ae34d4ac2c6aaa838ac917472ed23.diff

LOG: [CIR][AMDGPU] Add support for AMDGCN tanh builtins (#197852)

Adds codegen for the following AMDGCN tanh builtins:

- __builtin_amdgcn_tanhf (float)
- __builtin_amdgcn_tanhh (half)
- __builtin_amdgcn_tanh_bf16 (bfloat16)

These are lowered to the corresponding `llvm.amdgcn.tanh` intrinsic.

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
    clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index b681455c4e4b0..861ec451e2bc0 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -358,10 +358,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   case AMDGPU::BI__builtin_amdgcn_tanhf:
   case AMDGPU::BI__builtin_amdgcn_tanhh:
   case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
+    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.tanh").getValue();
   }
   case AMDGPU::BI__builtin_amdgcn_uicmp:
   case AMDGPU::BI__builtin_amdgcn_uicmpl:

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
index 1a1ab630e5a42..b6e46ba180c34 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -72,3 +72,27 @@ __device__ void test_exp2_bf16(__bf16* out, __bf16 a) {
 __device__ void test_log_bf16(__bf16* out, __bf16 a) {
   *out = __builtin_amdgcn_log_bf16(a);
 }
+
+// CIR-LABEL: @_Z14test_tanhf_f32Pff
+// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z14test_tanhf_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.tanh.f32(float %{{.*}})
+__device__ void test_tanhf_f32(float* out, float a) {
+  *out = __builtin_amdgcn_tanhf(a);
+}
+
+// CIR-LABEL: @_Z13test_tanh_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z13test_tanh_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.amdgcn.tanh.f16(half %{{.*}})
+__device__ void test_tanh_f16(_Float16* out, _Float16 a) {
+  *out = __builtin_amdgcn_tanhh(a);
+}
+
+// CIR-LABEL: @_Z14test_tanh_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.tanh" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z14test_tanh_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.tanh.bf16(bfloat %{{.*}})
+__device__ void test_tanh_bf16(__bf16* out, __bf16 a) {
+  *out = __builtin_amdgcn_tanh_bf16(a);
+}