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