From f1a00d80d96c7a91fca7ab0fdb115d61c05d05b6 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Mon, 9 Jun 2025 00:45:22 -0700 Subject: [PATCH 1/3] [SYCL][Devicelib] Add __devicelib_rint and fallback to __spirv_ocl_rint --- libdevice/cmath_wrapper.cpp | 9 +++++---- libdevice/cmath_wrapper_fp64.cpp | 9 +++++---- libdevice/device_math.h | 6 ++++++ libdevice/fallback-cmath-fp64.cpp | 4 ++++ libdevice/fallback-cmath.cpp | 3 +++ 5 files changed, 23 insertions(+), 8 deletions(-) diff --git a/libdevice/cmath_wrapper.cpp b/libdevice/cmath_wrapper.cpp index 1a5cec90d28cd..402759ce62199 100644 --- a/libdevice/cmath_wrapper.cpp +++ b/libdevice/cmath_wrapper.cpp @@ -198,9 +198,7 @@ float nearbyintf(float x) { return __nv_nearbyintf(x); } extern "C" SYCL_EXTERNAL float __nv_rintf(float); DEVICE_EXTERN_C_INLINE float rintf(float x) { return __nv_rintf(x); } -#endif // __NVPTX__ - -#ifdef __AMDGCN__ +#elif defined(__AMDGCN__) extern "C" SYCL_EXTERNAL float __ocml_nearbyint_f32(float); DEVICE_EXTERN_C_INLINE float nearbyintf(float x) { return __ocml_nearbyint_f32(x); } @@ -208,6 +206,9 @@ float nearbyintf(float x) { return __ocml_nearbyint_f32(x); } extern "C" SYCL_EXTERNAL float __ocml_rint_f32(float); DEVICE_EXTERN_C_INLINE float rintf(float x) { return __ocml_rint_f32(x); } -#endif // __AMDGCN__ +#else +DEVICE_EXTERN_C_INLINE +float rintf(float x) { return __devicelib_rintf(x); } +#endif #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ diff --git a/libdevice/cmath_wrapper_fp64.cpp b/libdevice/cmath_wrapper_fp64.cpp index 8d9c70ea0fe69..651a71d39594a 100644 --- a/libdevice/cmath_wrapper_fp64.cpp +++ b/libdevice/cmath_wrapper_fp64.cpp @@ -188,9 +188,7 @@ double nearbyint(double x) { return __nv_nearbyint(x); } extern "C" SYCL_EXTERNAL double __nv_rint(double); DEVICE_EXTERN_C_INLINE double rint(double x) { return __nv_rint(x); } -#endif // __NVPTX__ - -#ifdef __AMDGCN__ +#elif defined(__AMDGCN__) extern "C" SYCL_EXTERNAL double __ocml_nearbyint_f64(double); DEVICE_EXTERN_C_INLINE double nearbyint(double x) { return __ocml_nearbyint_f64(x); } @@ -198,7 +196,10 @@ double nearbyint(double x) { return __ocml_nearbyint_f64(x); } extern "C" SYCL_EXTERNAL double __ocml_rint_f64(double); DEVICE_EXTERN_C_INLINE double rint(double x) { return __ocml_rint_f64(x); } -#endif // __AMDGCN__ +#else +DEVICE_EXTERN_C_INLINE +double rint(double x) { return __devicelib_rint(x); } +#endif #if defined(_MSC_VER) #include diff --git a/libdevice/device_math.h b/libdevice/device_math.h index f4ee1711060c6..626728eb216c5 100644 --- a/libdevice/device_math.h +++ b/libdevice/device_math.h @@ -368,5 +368,11 @@ float __devicelib_scalbnf(float x, int n); DEVICE_EXTERN_C double __devicelib_scalbn(double x, int exp); +DEVICE_EXTERN_C +double __devicelib_rint(double x); + +DEVICE_EXTERN_C +float __devicelib_rintf(float x); + #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ #endif // __LIBDEVICE_DEVICE_MATH_H__ diff --git a/libdevice/fallback-cmath-fp64.cpp b/libdevice/fallback-cmath-fp64.cpp index d7e4364e2595d..03b563e2ae1e0 100644 --- a/libdevice/fallback-cmath-fp64.cpp +++ b/libdevice/fallback-cmath-fp64.cpp @@ -194,4 +194,8 @@ DEVICE_EXTERN_C_INLINE double __devicelib_scalbn(double x, int exp) { return __spirv_ocl_ldexp(x, exp); } + +DEVICE_EXTERN_C_INLINE +double __devicelib_rint(double x) { return __spirv_ocl_rint(x); } + #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ diff --git a/libdevice/fallback-cmath.cpp b/libdevice/fallback-cmath.cpp index 97cb4cf67b4c7..129a87aa2d597 100644 --- a/libdevice/fallback-cmath.cpp +++ b/libdevice/fallback-cmath.cpp @@ -215,4 +215,7 @@ float __devicelib_asinhf(float x) { return __spirv_ocl_asinh(x); } DEVICE_EXTERN_C_INLINE float __devicelib_atanhf(float x) { return __spirv_ocl_atanh(x); } +DEVICE_EXTERN_C_INLINE +float __devicelib_rintf(float x) { return __spirv_ocl_rint(x); } + #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ From b7b33853a20b02f592093f504d12b3de7efd8af9 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Mon, 9 Jun 2025 22:47:05 -0700 Subject: [PATCH 2/3] add test, remove fallback --- libdevice/cmath_wrapper.cpp | 2 +- libdevice/cmath_wrapper_fp64.cpp | 2 +- libdevice/device_math.h | 6 --- libdevice/fallback-cmath-fp64.cpp | 4 -- libdevice/fallback-cmath.cpp | 3 -- libdevice/test/check_cmath.txt | 12 ++++++ .../math-builtins/std_rint.cpp | 37 +++++++++++++++++++ 7 files changed, 51 insertions(+), 15 deletions(-) create mode 100644 libdevice/test/check_cmath.txt create mode 100644 sycl/test/check_device_code/math-builtins/std_rint.cpp diff --git a/libdevice/cmath_wrapper.cpp b/libdevice/cmath_wrapper.cpp index 402759ce62199..d59395b2d0994 100644 --- a/libdevice/cmath_wrapper.cpp +++ b/libdevice/cmath_wrapper.cpp @@ -208,7 +208,7 @@ DEVICE_EXTERN_C_INLINE float rintf(float x) { return __ocml_rint_f32(x); } #else DEVICE_EXTERN_C_INLINE -float rintf(float x) { return __devicelib_rintf(x); } +float rintf(float x) { return __spirv_ocl_rint(x); } #endif #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ diff --git a/libdevice/cmath_wrapper_fp64.cpp b/libdevice/cmath_wrapper_fp64.cpp index 651a71d39594a..720982799ea71 100644 --- a/libdevice/cmath_wrapper_fp64.cpp +++ b/libdevice/cmath_wrapper_fp64.cpp @@ -198,7 +198,7 @@ DEVICE_EXTERN_C_INLINE double rint(double x) { return __ocml_rint_f64(x); } #else DEVICE_EXTERN_C_INLINE -double rint(double x) { return __devicelib_rint(x); } +double rint(double x) { return __spirv_ocl_rint(x); } #endif #if defined(_MSC_VER) diff --git a/libdevice/device_math.h b/libdevice/device_math.h index 626728eb216c5..f4ee1711060c6 100644 --- a/libdevice/device_math.h +++ b/libdevice/device_math.h @@ -368,11 +368,5 @@ float __devicelib_scalbnf(float x, int n); DEVICE_EXTERN_C double __devicelib_scalbn(double x, int exp); -DEVICE_EXTERN_C -double __devicelib_rint(double x); - -DEVICE_EXTERN_C -float __devicelib_rintf(float x); - #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ #endif // __LIBDEVICE_DEVICE_MATH_H__ diff --git a/libdevice/fallback-cmath-fp64.cpp b/libdevice/fallback-cmath-fp64.cpp index 03b563e2ae1e0..d7e4364e2595d 100644 --- a/libdevice/fallback-cmath-fp64.cpp +++ b/libdevice/fallback-cmath-fp64.cpp @@ -194,8 +194,4 @@ DEVICE_EXTERN_C_INLINE double __devicelib_scalbn(double x, int exp) { return __spirv_ocl_ldexp(x, exp); } - -DEVICE_EXTERN_C_INLINE -double __devicelib_rint(double x) { return __spirv_ocl_rint(x); } - #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ diff --git a/libdevice/fallback-cmath.cpp b/libdevice/fallback-cmath.cpp index 129a87aa2d597..97cb4cf67b4c7 100644 --- a/libdevice/fallback-cmath.cpp +++ b/libdevice/fallback-cmath.cpp @@ -215,7 +215,4 @@ float __devicelib_asinhf(float x) { return __spirv_ocl_asinh(x); } DEVICE_EXTERN_C_INLINE float __devicelib_atanhf(float x) { return __spirv_ocl_atanh(x); } -DEVICE_EXTERN_C_INLINE -float __devicelib_rintf(float x) { return __spirv_ocl_rint(x); } - #endif // __SPIR__ || __SPIRV__ || __NVPTX__ || __AMDGCN__ diff --git a/libdevice/test/check_cmath.txt b/libdevice/test/check_cmath.txt new file mode 100644 index 0000000000000..856dea8c0e772 --- /dev/null +++ b/libdevice/test/check_cmath.txt @@ -0,0 +1,12 @@ +REQUIRES: libsycldevice + +Check functions in fp32 libdevice spirv file. + +RUN: llvm-spirv --spirv-target-env=SPV-IR -r %libsycldevice_spv_dir/libsycl-cmath.spv -o %t.bc +RUN: llvm-dis %t.bc -o %t.ll +RUN: FileCheck %s --input-file %t.ll + +CHECK: target triple ={{.*}}spir64 + +CHECK-LABEL: define spir_func float @rintf( +CHECK: call spir_func float @_Z16__spirv_ocl_rintf( diff --git a/sycl/test/check_device_code/math-builtins/std_rint.cpp b/sycl/test/check_device_code/math-builtins/std_rint.cpp new file mode 100644 index 0000000000000..435235a7c595e --- /dev/null +++ b/sycl/test/check_device_code/math-builtins/std_rint.cpp @@ -0,0 +1,37 @@ +// Make dump directory. +// RUN: rm -rf %t.spvdir && mkdir %t.spvdir + +// RUN: %clangxx -fsycl -fsycl-dump-device-code=%t.spvdir %s + +// Rename SPV file to explictly known filename. +// RUN: mv %t.spvdir/*.spv %t.spvdir/dump.spv + +// Convert to LLVM IR. +// RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spvdir/dump.spv +// RUN: llvm-dis %t.spvdir/dump.bc +// RUN: FileCheck --input-file=%t.spvdir/dump.ll %s + +#include +#include +#include + +// CHECK: call spir_func float @_Z16__spirv_ocl_rintf( + +using namespace sycl; + +int main() { + queue Q; + + float *Out = malloc_shared(1, Q); + Out[0] = 0.5f; + + try { + Q.submit([&](handler &Cgh) { + Cgh.parallel_for(nd_range<1>({1}, {1}), + [=](nd_item<1> Item) { *Out = std::rint(*Out); }); + }); + } catch (sycl::exception const &) { + } + + free(Out, Q); +} From f340020ddf22dab84ec2a1d568dd00bf6587d796 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Mon, 9 Jun 2025 23:13:13 -0700 Subject: [PATCH 3/3] remove sycl test, rint is lowered to llvm.rint.f32 intrinsic for spir target --- .../math-builtins/std_rint.cpp | 37 ------------------- 1 file changed, 37 deletions(-) delete mode 100644 sycl/test/check_device_code/math-builtins/std_rint.cpp diff --git a/sycl/test/check_device_code/math-builtins/std_rint.cpp b/sycl/test/check_device_code/math-builtins/std_rint.cpp deleted file mode 100644 index 435235a7c595e..0000000000000 --- a/sycl/test/check_device_code/math-builtins/std_rint.cpp +++ /dev/null @@ -1,37 +0,0 @@ -// Make dump directory. -// RUN: rm -rf %t.spvdir && mkdir %t.spvdir - -// RUN: %clangxx -fsycl -fsycl-dump-device-code=%t.spvdir %s - -// Rename SPV file to explictly known filename. -// RUN: mv %t.spvdir/*.spv %t.spvdir/dump.spv - -// Convert to LLVM IR. -// RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spvdir/dump.spv -// RUN: llvm-dis %t.spvdir/dump.bc -// RUN: FileCheck --input-file=%t.spvdir/dump.ll %s - -#include -#include -#include - -// CHECK: call spir_func float @_Z16__spirv_ocl_rintf( - -using namespace sycl; - -int main() { - queue Q; - - float *Out = malloc_shared(1, Q); - Out[0] = 0.5f; - - try { - Q.submit([&](handler &Cgh) { - Cgh.parallel_for(nd_range<1>({1}, {1}), - [=](nd_item<1> Item) { *Out = std::rint(*Out); }); - }); - } catch (sycl::exception const &) { - } - - free(Out, Q); -}