Skip to content

clang/HIP: Use generic builtins for f32 exp and log #129638

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Mar 6, 2025
Merged

clang/HIP: Use generic builtins for f32 exp and log #129638

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions clang/lib/Headers/__clang_hip_math.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -392,7 +392,7 @@ __DEVICE__
float erfinvf(float __x) { return __ocml_erfinv_f32(__x); }

__DEVICE__
float exp10f(float __x) { return __ocml_exp10_f32(__x); }
float exp10f(float __x) { return __builtin_exp10f(__x); }

__DEVICE__
float exp2f(float __x) { return __builtin_exp2f(__x); }
Expand Down Expand Up @@ -495,13 +495,13 @@ __DEVICE__
float log1pf(float __x) { return __ocml_log1p_f32(__x); }

__DEVICE__
float log2f(float __x) { return __FAST_OR_SLOW(__log2f, __ocml_log2_f32)(__x); }
float log2f(float __x) { return __FAST_OR_SLOW(__log2f, __builtin_log2f)(__x); }

__DEVICE__
float logbf(float __x) { return __ocml_logb_f32(__x); }

__DEVICE__
float logf(float __x) { return __FAST_OR_SLOW(__logf, __ocml_log_f32)(__x); }
float logf(float __x) { return __FAST_OR_SLOW(__logf, __builtin_logf)(__x); }

__DEVICE__
long int lrintf(float __x) { return __builtin_rintf(__x); }
Expand Down
42 changes: 20 additions & 22 deletions clang/test/Headers/__clang_hip_math.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1075,7 +1075,6 @@ extern "C" __device__ double test_cospi(double x) {
return cospi(x);
}

//
// DEFAULT-LABEL: @test_cyl_bessel_i0f(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_i0_f32(float noundef [[X:%.*]]) #[[ATTR14]]
Expand Down Expand Up @@ -1270,23 +1269,23 @@ extern "C" __device__ double test_erfinv(double x) {

// DEFAULT-LABEL: @test_exp10f(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// DEFAULT-NEXT: ret float [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp10.f32(float [[X:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_exp10f(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_exp10_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.exp10.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_exp10f(
// APPROX-NEXT: entry:
// APPROX-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// APPROX-NEXT: ret float [[CALL_I]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.exp10.f32(float [[X:%.*]])
// APPROX-NEXT: ret float [[TMP0]]
//
// AMDGCNSPIRV-LABEL: @test_exp10f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_exp10_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// AMDGCNSPIRV-NEXT: ret float [[CALL_I]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.exp10.f32(float [[X:%.*]])
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_exp10f(float x) {
return exp10f(x);
Expand Down Expand Up @@ -1748,7 +1747,6 @@ extern "C" __device__ double test_fmax(double x, double y) {
return fmax(x, y);
}

//
// DEFAULT-LABEL: @test_fminf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
Expand Down Expand Up @@ -2823,13 +2821,13 @@ extern "C" __device__ double test_log1p(double x) {

// DEFAULT-LABEL: @test_log2f(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// DEFAULT-NEXT: ret float [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log2.f32(float [[X:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_log2f(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_log2_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log2.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_log2f(
// APPROX-NEXT: entry:
Expand All @@ -2838,8 +2836,8 @@ extern "C" __device__ double test_log1p(double x) {
//
// AMDGCNSPIRV-LABEL: @test_log2f(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_log2_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// AMDGCNSPIRV-NEXT: ret float [[CALL_I]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log2.f32(float [[X:%.*]])
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_log2f(float x) {
return log2f(x);
Expand Down Expand Up @@ -2919,13 +2917,13 @@ extern "C" __device__ double test_logb(double x) {

// DEFAULT-LABEL: @test_logf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract noundef float @__ocml_log_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// DEFAULT-NEXT: ret float [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_logf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @__ocml_log_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR13]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: @test_logf(
// APPROX-NEXT: entry:
Expand All @@ -2934,8 +2932,8 @@ extern "C" __device__ double test_logb(double x) {
//
// AMDGCNSPIRV-LABEL: @test_logf(
// AMDGCNSPIRV-NEXT: entry:
// AMDGCNSPIRV-NEXT: [[CALL_I:%.*]] = tail call contract spir_func noundef addrspace(4) float @__ocml_log_f32(float noundef [[X:%.*]]) #[[ATTR13]]
// AMDGCNSPIRV-NEXT: ret float [[CALL_I]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log.f32(float [[X:%.*]])
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_logf(float x) {
return logf(x);
Expand Down