[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