[clang] 26c200d - [CIR][AMDGPU] Add support for AMDGCN sin builtins (#197784)

[via cfe-commits]

Author: Ayokunle Amodu
Date: 2026-07-16T09:53:32-04:00
New Revision: 26c200d88e720fa88e8692ace13a493696ac591f

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

LOG: [CIR][AMDGPU] Add support for AMDGCN sin builtins (#197784)

Adds codegen for the following AMDGCN sin builtins:

- __builtin_amdgcn_sinf (float)
- __builtin_amdgcn_sinh (half)
- __builtin_amdgcn_sin_bf16 (bfloat16)

These are lowered to the corresponding `llvm.amdgcn.sin` 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 0b60faeab4d04..2f0ef3a50bc86 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -273,10 +273,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   case AMDGPU::BI__builtin_amdgcn_sinf:
   case AMDGPU::BI__builtin_amdgcn_sinh:
   case AMDGPU::BI__builtin_amdgcn_sin_bf16: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
+    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.sin").getValue();
   }
   case AMDGPU::BI__builtin_amdgcn_cosf:
   case AMDGPU::BI__builtin_amdgcn_cosh:

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
index b06f5a666d54e..2fc4daf2a29ab 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -40,3 +40,11 @@ __device__ void test_rsq_bf16(__bf16* out, __bf16 a) {
 __device__ void test_sqrt_bf16(__bf16* out, __bf16 a) {
   *out = __builtin_amdgcn_sqrt_bf16(a);
 }
+
+// CIR-LABEL: @_Z13test_sin_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z13test_sin_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.sin.bf16(bfloat %{{.*}})
+__device__ void test_sin_bf16(__bf16* out, __bf16 a) {
+  *out = __builtin_amdgcn_sin_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 6d4d14285cb57..743b0e15c2f24 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -87,3 +87,11 @@ __device__ void test_rsq_h(_Float16* out, _Float16 a) {
 __device__ void test_sqrt_f16(_Float16* out, _Float16 a) {
   *out = __builtin_amdgcn_sqrth(a);
 }
+
+// CIR-LABEL: @_Z12test_sin_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z12test_sin_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.amdgcn.sin.f16(half %{{.*}})
+__device__ void test_sin_f16(_Float16* out, _Float16 a) {
+  *out = __builtin_amdgcn_sinh(a);
+}

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index dcb155a66270d..cf65e277beabe 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -175,3 +175,12 @@ __device__ void test_sqrt_f32(float* out, float a) {
 __device__ void test_sqrt_f64(double* out, double a) {
   *out = __builtin_amdgcn_sqrt(a);
 }
+
+// CIR-LABEL: @_Z13test_sinf_f32Pff
+// CIR: cir.call_llvm_intrinsic "amdgcn.sin" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z13test_sinf_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.sin.f32(float %{{.*}})
+__device__ void test_sinf_f32(float* out, float a) {
+  *out = __builtin_amdgcn_sinf(a);
+}
+