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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
