[PATCH] D154790: [HIP] Use native math functions for `-fcuda-approx-transcendentals`
Yaxun Liu via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Wed Jul 12 13:32:25 PDT 2023
yaxunl marked 8 inline comments as done.
yaxunl added inline comments.
================
Comment at: clang/lib/Headers/__clang_hip_math.h:160
+__DEVICE__
+float __exp10f(float __x) { return __ocml_native_exp10_f32(__x); }
+
----------------
arsenm wrote:
> We should have llvm.exp10 but don't today. Just inline __builtin_exp2f(M_LOG2_10_F * x)
done
================
Comment at: clang/lib/Headers/__clang_hip_math.h:163
+__DEVICE__
+float __expf(float __x) { return __ocml_native_exp_f32(__x); }
+
----------------
arsenm wrote:
> arsenm wrote:
> > __builtin_expf
> Maybe this should just be __builtin_amdgcn_exp2f
__builtin_expf causes extra calculations. will use __builtin_amdgcn_exp2f
================
Comment at: clang/lib/Headers/__clang_hip_math.h:280
+__DEVICE__
+float __log10f(float __x) { return __ocml_native_log10_f32(__x); }
+
----------------
arsenm wrote:
> 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
will use __builtin_log10f
================
Comment at: clang/lib/Headers/__clang_hip_math.h:283
+__DEVICE__
+float __log2f(float __x) { return __ocml_native_log2_f32(__x); }
+
----------------
arsenm wrote:
> __builtin_log2f or __builtin_amdgcn_log. Ideally would be llvm.log2.f32 with afn set
will use `__builtin_amdgcn_logf`
================
Comment at: clang/lib/Headers/__clang_hip_math.h:286
+__DEVICE__
+float __logf(float __x) { return __ocml_native_log_f32(__x); }
+
----------------
arsenm wrote:
> Same as log10 case, except with log
will use `__builtin_logf`
================
Comment at: clang/lib/Headers/__clang_hip_math.h:304
+__DEVICE__
+float __tanf(float __x) { return __ocml_tan_f32(__x); }
+// END INTRINSICS
----------------
b-sumner wrote:
> We could consider multiplying native_sin here with the native_recip of native_cos.
will do
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D154790/new/
https://reviews.llvm.org/D154790
More information about the cfe-commits
mailing list