[clang] dfa1360 - [CIR][AMDGPU] Add support for AMDGCN cos builtins (#197806)

[via cfe-commits]

Author: Ayokunle Amodu
Date: 2026-07-16T11:42:39-04:00
New Revision: dfa1360735f5a834ce96c8c63594e2cdc3bfe0ea

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

LOG: [CIR][AMDGPU] Add support for AMDGCN cos builtins (#197806)

Adds codegen for the following AMDGCN cos builtins:

- __builtin_amdgcn_cosf (float)
- __builtin_amdgcn_cosh (half)
- __builtin_amdgcn_cos_bf16 (bfloat16)

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

Added: 
    

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
    clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
    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 2f0ef3a50bc86..cdb7106b84094 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -278,10 +278,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   case AMDGPU::BI__builtin_amdgcn_cosf:
   case AMDGPU::BI__builtin_amdgcn_cosh:
   case AMDGPU::BI__builtin_amdgcn_cos_bf16: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
+    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.cos").getValue();
   }
   case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
     return emitAMDGPUDispatchPtr(*this, expr);

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
index 2fc4daf2a29ab..b85141d99d783 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -48,3 +48,11 @@ __device__ void test_sqrt_bf16(__bf16* out, __bf16 a) {
 __device__ void test_sin_bf16(__bf16* out, __bf16 a) {
   *out = __builtin_amdgcn_sin_bf16(a);
 }
+
+// CIR-LABEL: @_Z13test_cos_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.cos" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z13test_cos_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.cos.bf16(bfloat %{{.*}})
+__device__ void test_cos_bf16(__bf16* out, __bf16 a) {
+  *out = __builtin_amdgcn_cos_bf16(a);
+}

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
index 743b0e15c2f24..ce5168996ab28 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -95,3 +95,11 @@ __device__ void test_sqrt_f16(_Float16* out, _Float16 a) {
 __device__ void test_sin_f16(_Float16* out, _Float16 a) {
   *out = __builtin_amdgcn_sinh(a);
 }
+
+// CIR-LABEL: @_Z12test_cos_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.cos" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z12test_cos_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.amdgcn.cos.f16(half %{{.*}})
+__device__ void test_cos_f16(_Float16* out, _Float16 a) {
+  *out = __builtin_amdgcn_cosh(a);
+}

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index cf65e277beabe..2713f71a3cdf2 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -184,3 +184,10 @@ __device__ void test_sinf_f32(float* out, float a) {
   *out = __builtin_amdgcn_sinf(a);
 }
 
+// CIR-LABEL: @_Z13test_cosf_f32Pff
+// CIR: cir.call_llvm_intrinsic "amdgcn.cos" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z13test_cosf_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.cos.f32(float %{{.*}})
+__device__ void test_cosf_f32(float* out, float a) {
+  *out = __builtin_amdgcn_cosf(a);
+}