Author: Ayokunle Amodu Date: 2026-07-13T06:21:15-04:00 New Revision: da138053e78c5af7b83f49674b4e190e1803102d
URL: https://github.com/llvm/llvm-project/commit/da138053e78c5af7b83f49674b4e190e1803102d DIFF: https://github.com/llvm/llvm-project/commit/da138053e78c5af7b83f49674b4e190e1803102d.diff LOG: [CIR][AMDGPU] Add support for AMDGCN rcp builtins (#197447) Adds codegen for the following AMDGCN reciprocal builtins: - __builtin_amdgcn_rcp (double) - __builtin_amdgcn_rcpf (float) - __builtin_amdgcn_rcph (half) - __builtin_amdgcn_rcp_bf16 (bfloat16) These are lowered to the corresponding `llvm.amdgcn.rcp` intrinsic. Added: clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp 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 0a7ba0c194400..5736eb70bc123 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -251,10 +251,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_rcpf: case AMDGPU::BI__builtin_amdgcn_rcph: case AMDGPU::BI__builtin_amdgcn_rcp_bf16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rcp").getValue(); } case AMDGPU::BI__builtin_amdgcn_sqrt: case AMDGPU::BI__builtin_amdgcn_sqrtf: diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip new file mode 100644 index 0000000000000..64fa56b45e39a --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -0,0 +1,26 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +//===----------------------------------------------------------------------===// +// Test AMDGPU builtins +//===----------------------------------------------------------------------===// + +#define __device__ __attribute__((device)) + +// CIR-LABEL: @_Z13test_rcp_bf16PDF16bDF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.bf16) -> !cir.bf16 +// LLVM: define{{.*}} void @_Z13test_rcp_bf16PDF16bDF16b +// LLVM: call{{.*}} bfloat @llvm.amdgcn.rcp.bf16(bfloat %{{.*}}) +__device__ void test_rcp_bf16(__bf16* out, __bf16 a) { + *out = __builtin_amdgcn_rcp_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 c21359ff7a88b..3ee9f1664f61a 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip @@ -63,3 +63,11 @@ __device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) { *out = __builtin_amdgcn_div_fixuph(a, b, c); } + +// CIR-LABEL: @_Z12test_rcp_f16PDF16_DF16_ +// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.f16) -> !cir.f16 +// LLVM: define{{.*}} void @_Z12test_rcp_f16PDF16_DF16_ +// LLVM: call{{.*}} half @llvm.amdgcn.rcp.f16(half %{{.*}}) +__device__ void test_rcp_f16(_Float16* out, _Float16 a) { + *out = __builtin_amdgcn_rcph(a); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index ca708bca8587b..9e279d66de1fa 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -111,3 +111,19 @@ __device__ void test_readfirstlane(int* out, int a) { __device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out) { *out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr(); } + +// CIR-LABEL: @_Z12test_rcp_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z12test_rcp_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.rcp.f32(float %{{.*}}) +__device__ void test_rcp_f32(float* out, float a) { + *out = __builtin_amdgcn_rcpf(a); +} + +// CIR-LABEL: @_Z12test_rcp_f64Pdd +// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.double) -> !cir.double +// LLVM: define{{.*}} void @_Z12test_rcp_f64Pdd +// LLVM: call{{.*}} double @llvm.amdgcn.rcp.f64(double %{{.*}}) +__device__ void test_rcp_f64(double* out, double a) { + *out = __builtin_amdgcn_rcp(a); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
