[PATCH] D154790: [HIP] Use native math functions for `-fcuda-approx-transcendentals`

Matt Arsenault via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Sun Jul 9 16:22:18 PDT 2023


arsenm added a comment.

This would be a lot easier if we clang fp was fully featured. As far as I can tell it lets you set #pragma clang fp reassociate(on), and contract(fast), but doesn't have a way to set arcp, afn, ninf or nnan



================
Comment at: clang/lib/Headers/__clang_hip_math.h:160
+__DEVICE__
+float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); }
+
----------------
We should have llvm.exp10 but don't today. Just inline __builtin_exp2f(M_LOG2_10_F * x)


================
Comment at: clang/lib/Headers/__clang_hip_math.h:163
+__DEVICE__
+float __expf(float __x) { return __ocml_native_exp_f32(__x); }
+
----------------
__builtin_expf 


================
Comment at: clang/lib/Headers/__clang_hip_math.h:280
+__DEVICE__
+float __log10f(float __x) { return __ocml_native_log10_f32(__x); }
+
----------------
We want llvm.log10.f32 with afn set. You get closer by just using __builtin_log10f. Ideally we would have a pragma to set afn locally 


================
Comment at: clang/lib/Headers/__clang_hip_math.h:283
+__DEVICE__
+float __log2f(float __x) { return __ocml_native_log2_f32(__x); }
+
----------------
__builtin_log2f or __builtin_amdgcn_log. Ideally would be llvm.log2.f32 with afn set 


================
Comment at: clang/lib/Headers/__clang_hip_math.h:286
+__DEVICE__
+float __logf(float __x) { return __ocml_native_log_f32(__x); }
+
----------------
Same as log10 case, except with log 


================
Comment at: clang/lib/Headers/__clang_hip_math.h:301
+__DEVICE__
+float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
+
----------------
Really we want these inlined with FMF set, the built libraries currently have no flags 


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

https://reviews.llvm.org/D154790



More information about the cfe-commits mailing list