https://github.com/robertvirany updated https://github.com/llvm/llvm-project/pull/205661
>From f25cb916f2bc6e5ab3b62eb5a9420a09a60cae20 Mon Sep 17 00:00:00 2001 From: Robert Virany <[email protected]> Date: Wed, 24 Jun 2026 15:49:37 -0400 Subject: [PATCH 1/2] [CUDA] Lower device sqrtf through builtin sqrt Lower CUDA device sqrtf through __builtin_sqrtf instead of the libdevice __nv_sqrtf wrapper. This lets NVPTX select precise or approximate sqrt from the LLVM sqrt intrinsic and fast-math flags. Fixes #131749 --- clang/lib/Headers/__clang_cuda_math.h | 2 +- clang/test/CodeGenCUDA/sqrtf-precise.cu | 27 +++++++++++++++++++++++++ 2 files changed, 28 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenCUDA/sqrtf-precise.cu diff --git a/clang/lib/Headers/__clang_cuda_math.h b/clang/lib/Headers/__clang_cuda_math.h index 44c6e9a4e48d1..cf361e6109194 100644 --- a/clang/lib/Headers/__clang_cuda_math.h +++ b/clang/lib/Headers/__clang_cuda_math.h @@ -315,7 +315,7 @@ __DEVICE__ float sinhf(float __a) { return __nv_sinhf(__a); } __DEVICE__ double sinpi(double __a) { return __nv_sinpi(__a); } __DEVICE__ float sinpif(float __a) { return __nv_sinpif(__a); } __DEVICE__ double sqrt(double __a) { return __nv_sqrt(__a); } -__DEVICE__ float sqrtf(float __a) { return __nv_sqrtf(__a); } +__DEVICE__ float sqrtf(float __a) { return __builtin_sqrtf(__a); } __DEVICE__ double tan(double __a) { return __nv_tan(__a); } __DEVICE__ float tanf(float __a) { return __nv_tanf(__a); } __DEVICE__ double tanh(double __a) { return __nv_tanh(__a); } diff --git a/clang/test/CodeGenCUDA/sqrtf-precise.cu b/clang/test/CodeGenCUDA/sqrtf-precise.cu new file mode 100644 index 0000000000000..5ba72c376b71a --- /dev/null +++ b/clang/test/CodeGenCUDA/sqrtf-precise.cu @@ -0,0 +1,27 @@ +// REQUIRES: nvptx-registered-target + +// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \ +// RUN: -target-cpu sm_75 -O2 -S -o - %s \ +// RUN: -include __clang_cuda_runtime_wrapper.h \ +// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ +// RUN: -internal-isystem %S/../Headers/Inputs/include \ +// RUN: | FileCheck --check-prefix=RN %s + +// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \ +// RUN: -target-cpu sm_75 -O2 -fapprox-func -S -o - %s \ +// RUN: -include __clang_cuda_runtime_wrapper.h \ +// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ +// RUN: -internal-isystem %S/../Headers/Inputs/include \ +// RUN: | FileCheck --check-prefix=APPROX %s + +#include <math.h> + +// RN-LABEL: .func{{.*}}_Z1ff +// RN: sqrt.rn.f32 +// RN-NOT: sqrt.approx.f32 +// +// APPROX-LABEL: .func{{.*}}_Z1ff +// APPROX: sqrt.approx.f32 +__device__ float f(float x) { + return sqrtf(x); +} >From 020600ffb3829b5e62d46303c44ddb356dcdeae5 Mon Sep 17 00:00:00 2001 From: Robert Virany <[email protected]> Date: Thu, 16 Jul 2026 10:42:23 -0600 Subject: [PATCH 2/2] [CUDA] Make sqrtf builtin test self-contained --- clang/test/CodeGenCUDA/sqrtf-precise.cu | 38 ++++++++++++++++--------- 1 file changed, 24 insertions(+), 14 deletions(-) diff --git a/clang/test/CodeGenCUDA/sqrtf-precise.cu b/clang/test/CodeGenCUDA/sqrtf-precise.cu index 5ba72c376b71a..bcef140ef773c 100644 --- a/clang/test/CodeGenCUDA/sqrtf-precise.cu +++ b/clang/test/CodeGenCUDA/sqrtf-precise.cu @@ -2,26 +2,36 @@ // RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \ // RUN: -target-cpu sm_75 -O2 -S -o - %s \ -// RUN: -include __clang_cuda_runtime_wrapper.h \ -// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ -// RUN: -internal-isystem %S/../Headers/Inputs/include \ -// RUN: | FileCheck --check-prefix=RN %s +// RUN: | FileCheck --check-prefix=DEFAULT %s // RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \ // RUN: -target-cpu sm_75 -O2 -fapprox-func -S -o - %s \ -// RUN: -include __clang_cuda_runtime_wrapper.h \ -// RUN: -internal-isystem %S/../../lib/Headers/cuda_wrappers \ -// RUN: -internal-isystem %S/../Headers/Inputs/include \ // RUN: | FileCheck --check-prefix=APPROX %s -#include <math.h> +// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \ +// RUN: -target-cpu sm_75 -O2 -fapprox-func \ +// RUN: -mllvm -nvptx-prec-sqrtf32=1 -S -o - %s \ +// RUN: | FileCheck --check-prefix=FORCE-PRECISE %s + +// RUN: %clang_cc1 -fcuda-is-device -triple nvptx64-nvidia-cuda \ +// RUN: -target-cpu sm_75 -O2 -mllvm -nvptx-prec-sqrtf32=0 -S -o - %s \ +// RUN: | FileCheck --check-prefix=FORCE-APPROX %s -// RN-LABEL: .func{{.*}}_Z1ff -// RN: sqrt.rn.f32 -// RN-NOT: sqrt.approx.f32 +// DEFAULT-LABEL: .func{{.*}}builtin_sqrtf +// DEFAULT: sqrt.rn.f32 +// DEFAULT-NOT: sqrt.approx.f32 // -// APPROX-LABEL: .func{{.*}}_Z1ff +// APPROX-LABEL: .func{{.*}}builtin_sqrtf // APPROX: sqrt.approx.f32 -__device__ float f(float x) { - return sqrtf(x); +// APPROX-NOT: sqrt.rn.f32 +// +// FORCE-PRECISE-LABEL: .func{{.*}}builtin_sqrtf +// FORCE-PRECISE: sqrt.rn.f32 +// FORCE-PRECISE-NOT: sqrt.approx.f32 +// +// FORCE-APPROX-LABEL: .func{{.*}}builtin_sqrtf +// FORCE-APPROX: sqrt.approx.f32 +// FORCE-APPROX-NOT: sqrt.rn.f32 +extern "C" __attribute__((device)) float builtin_sqrtf(float x) { + return __builtin_sqrtf(x); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
