[libc-commits] [PATCH] D152603: [libc] Add math functions to AMD/NVPTX libm

[Matt Arsenault via Phabricator via libc-commits]

arsenm added inline comments.


================
Comment at: libc/src/math/gpu/frexp.cpp:19-28
+LLVM_LIBC_FUNCTION(double, frexp, (double x, int *p)) {
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+  return __nv_frexp(x, p);
+#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+  *p = __builtin_amdgcn_frexp_exp(x);
+  return __builtin_amdgcn_frexp_mant(x);
+#else
----------------
See https://reviews.llvm.org/D149715, this should get a generic builtin


================
Comment at: libc/src/math/gpu/frexpf.cpp:19
+
+LLVM_LIBC_FUNCTION(float, frexpf, (float x, int *p)) {
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
----------------
Same, https://reviews.llvm.org/D149715

Also, I think splitting the files up by FP type is pretty annoying. Many of the functions would benefit from templating


================
Comment at: libc/src/math/gpu/vendor/amdgpu/amdgpu.h:28
+LIBC_INLINE float coshf(float x) { return __ocml_cosh_f32(x); }
+LIBC_INLINE float expf(float x) { return __ocml_exp_f32(x); }
+LIBC_INLINE float exp2f(float x) { return __ocml_exp2_f32(x); }
----------------
Will be correctly lowered and should use the builtin after D153025


================
Comment at: libc/src/math/gpu/vendor/amdgpu/amdgpu.h:29
+LIBC_INLINE float expf(float x) { return __ocml_exp_f32(x); }
+LIBC_INLINE float exp2f(float x) { return __ocml_exp2_f32(x); }
+LIBC_INLINE float exp10f(float x) { return __ocml_exp10_f32(x); }
----------------
Will be correctly lowered after D153024 and should use the builtin


================
Comment at: libc/src/math/gpu/vendor/amdgpu/amdgpu.h:30
+LIBC_INLINE float exp2f(float x) { return __ocml_exp2_f32(x); }
+LIBC_INLINE float exp10f(float x) { return __ocml_exp10_f32(x); }
+LIBC_INLINE float expm1f(float x) { return __ocml_expm1_f32(x); }
----------------
In the abstract sense covered by D153024 but there's currently no defined llvm.exp10. This is asymmetrical and annoying and should be fixed, then the builtin will work


================
Comment at: libc/src/math/gpu/vendor/amdgpu/amdgpu.h:38-39
+LIBC_INLINE int ilogbf(float x) { return __ocml_ilogb_f32(x); }
+LIBC_INLINE double ldexp(double x, int i) { return __ocml_ldexp_f64(x, i); }
+LIBC_INLINE float ldexpf(float x, int i) { return __ocml_ldexp_f32(x, i); }
+LIBC_INLINE long long llrint(double x) { return __ocml_rint_f64(x); }
----------------
Should use the new __builtin_ldexp/__builtin_ldexpf


================
Comment at: libc/src/math/gpu/vendor/amdgpu/amdgpu.h:40-41
+LIBC_INLINE float ldexpf(float x, int i) { return __ocml_ldexp_f32(x, i); }
+LIBC_INLINE long long llrint(double x) { return __ocml_rint_f64(x); }
+LIBC_INLINE long long llrintf(float x) { return __ocml_rint_f32(x); }
+LIBC_INLINE long long llround(double x) { return __ocml_round_f64(x); }
----------------
__builtin_rint


================
Comment at: libc/src/math/gpu/vendor/amdgpu/amdgpu.h:42-43
+LIBC_INLINE long long llrintf(float x) { return __ocml_rint_f32(x); }
+LIBC_INLINE long long llround(double x) { return __ocml_round_f64(x); }
+LIBC_INLINE long long llroundf(float x) { return __ocml_round_f32(x); }
 LIBC_INLINE double sin(double x) { return __ocml_sin_f64(x); }
----------------
__builtin_round


================
Comment at: libc/src/math/gpu/vendor/amdgpu/amdgpu.h:45
 LIBC_INLINE double sin(double x) { return __ocml_sin_f64(x); }
 
 } // namespace internal
----------------
pow/powf?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D152603/new/

https://reviews.llvm.org/D152603