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