[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