Author: Ayokunle Amodu Date: 2026-07-15T08:46:37-04:00 New Revision: 9744b475e28e180599fd1a42d956f74b0a1aaed8
URL: https://github.com/llvm/llvm-project/commit/9744b475e28e180599fd1a42d956f74b0a1aaed8 DIFF: https://github.com/llvm/llvm-project/commit/9744b475e28e180599fd1a42d956f74b0a1aaed8.diff LOG: [CIR][AMDGPU] Add support for AMDGCN rsq and rsq_clamp builtins (#197349) Adds codegen for the following AMDGCN reciprocal square root builtins: - __builtin_amdgcn_rsq (double) - __builtin_amdgcn_rsqf (float) - __builtin_amdgcn_rsqh (half) - __builtin_amdgcn_rsq_bf16 (bfloat16) - __builtin_amdgcn_rsq_clamp (double) - __builtin_amdgcn_rsq_clampf (float) These are lowered to the corresponding `llvm.amdgcn.rsq` and `llvm.amdgcn.rsq.clamp` intrinsic calls. 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 5736eb70bc123..724171fd2fa0e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -266,17 +266,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, case AMDGPU::BI__builtin_amdgcn_rsqf: case AMDGPU::BI__builtin_amdgcn_rsqh: case AMDGPU::BI__builtin_amdgcn_rsq_bf16: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq").getValue(); } case AMDGPU::BI__builtin_amdgcn_rsq_clamp: case AMDGPU::BI__builtin_amdgcn_rsq_clampf: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rsq.clamp") + .getValue(); } case AMDGPU::BI__builtin_amdgcn_sinf: case AMDGPU::BI__builtin_amdgcn_sinh: diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip index 64fa56b45e39a..4a5a2a207076f 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip @@ -11,12 +11,12 @@ // RUN: -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s +#define __device__ __attribute__((device)) + //===----------------------------------------------------------------------===// // 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 @@ -24,3 +24,11 @@ __device__ void test_rcp_bf16(__bf16* out, __bf16 a) { *out = __builtin_amdgcn_rcp_bf16(a); } + +// CIR-LABEL: @_Z13test_rsq_bf16PDF16bDF16b +// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.bf16) -> !cir.bf16 +// LLVM: define{{.*}} void @_Z13test_rsq_bf16PDF16bDF16b +// LLVM: call{{.*}} bfloat @llvm.amdgcn.rsq.bf16(bfloat %{{.*}}) +__device__ void test_rsq_bf16(__bf16* out, __bf16 a) { + *out = __builtin_amdgcn_rsq_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 3ee9f1664f61a..d0c6f89842730 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip @@ -71,3 +71,11 @@ __device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float __device__ void test_rcp_f16(_Float16* out, _Float16 a) { *out = __builtin_amdgcn_rcph(a); } + +// CIR-LABEL: @_Z10test_rsq_hPDF16_DF16_ +// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.f16) -> !cir.f16 +// LLVM: define{{.*}} void @_Z10test_rsq_hPDF16_DF16_ +// LLVM: call{{.*}} half @llvm.amdgcn.rsq.f16(half %{{.*}}) +__device__ void test_rsq_h(_Float16* out, _Float16 a) { + *out = __builtin_amdgcn_rsqh(a); +} diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip index 9e279d66de1fa..7114aec6864c2 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -127,3 +127,35 @@ __device__ void test_rcp_f32(float* out, float a) { __device__ void test_rcp_f64(double* out, double a) { *out = __builtin_amdgcn_rcp(a); } + +// CIR-LABEL: @_Z12test_rsq_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z12test_rsq_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.rsq.f32(float %{{.*}}) +__device__ void test_rsq_f32(float* out, float a) { + *out = __builtin_amdgcn_rsqf(a); +} + +// CIR-LABEL: @_Z12test_rsq_f64Pdd +// CIR: cir.call_llvm_intrinsic "amdgcn.rsq" {{.*}} : (!cir.double) -> !cir.double +// LLVM: define{{.*}} void @_Z12test_rsq_f64Pdd +// LLVM: call{{.*}} double @llvm.amdgcn.rsq.f64(double %{{.*}}) +__device__ void test_rsq_f64(double* out, double a) { + *out = __builtin_amdgcn_rsq(a); +} + +// CIR-LABEL: @_Z18test_rsq_clamp_f32Pff +// CIR: cir.call_llvm_intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.float) -> !cir.float +// LLVM: define{{.*}} void @_Z18test_rsq_clamp_f32Pff +// LLVM: call{{.*}} float @llvm.amdgcn.rsq.clamp.f32(float %{{.*}}) +__device__ void test_rsq_clamp_f32(float* out, float a) { + *out = __builtin_amdgcn_rsq_clampf(a); +} + +// CIR-LABEL: @_Z18test_rsq_clamp_f64Pdd +// CIR: cir.call_llvm_intrinsic "amdgcn.rsq.clamp" {{.*}} : (!cir.double) -> !cir.double +// LLVM: define{{.*}} void @_Z18test_rsq_clamp_f64Pdd +// LLVM: call{{.*}} double @llvm.amdgcn.rsq.clamp.f64(double %{{.*}}) +__device__ void test_rsq_clamp_f64(double* out, double a) { + *out = __builtin_amdgcn_rsq_clamp(a); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
