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

Jump to bottom
Merged
arsenm merged 1 commit into from
Mar 6, 2025
Merged

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

arsenm merged 1 commit into from
Mar 6, 2025

Conversation

@arsenm
Copy link
Contributor

@arsenm arsenm commented Mar 4, 2025

Stop using ocml declarations which are just a shim around the
intrinsics.

@arsenm
clang/HIP: Use generic builtins for f32 exp and log
0b61a5f 
Stop using ocml declarations which are just a shim around the
intrinsics.
@arsenm Graphite App
Copy link
Contributor Author

arsenm commented Mar 4, 2025 •
edited
Loading

This stack of pull requests is managed by Graphite. Learn more about stacking.

@arsenm arsenm added the clang:headers Headers provided by Clang, e.g. for intrinsics label Mar 4, 2025 — with Graphite App
@arsenm arsenm requested review from AlexVlx and yxsamliu March 4, 2025 03:56
@arsenm arsenm marked this pull request as ready for review March 4, 2025 03:57
@llvmbot llvmbot added clang Clang issues not falling into any other category backend:X86 labels Mar 4, 2025
@llvmbot
Copy link
Member

llvmbot commented Mar 4, 2025 •
edited
Loading

@llvm/pr-subscribers-backend-x86

@llvm/pr-subscribers-clang

Author: Matt Arsenault (arsenm)

Changes

Stop using ocml declarations which are just a shim around the
intrinsics.

Full diff: https://github.com/llvm/llvm-project/pull/129638.diff

2 Files Affected:

  • (modified) clang/lib/Headers/__clang_hip_math.h (+3-3)
  • (modified) clang/test/Headers/__clang_hip_math.hip (+20-22)
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
index bf8517bc3a507..51d9acbb87270 100644
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -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); }
@@ -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); }
diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip
index d448ab134ca4d..ff9f55a8e0710 100644
--- a/clang/test/Headers/__clang_hip_math.hip
+++ b/clang/test/Headers/__clang_hip_math.hip
@@ -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]]
@@ -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);
@@ -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:%.*]])
@@ -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:
@@ -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);
@@ -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:
@@ -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);

@arsenm arsenm mentioned this pull request Mar 4, 2025
Merged
AlexVlx
AlexVlx approved these changes Mar 4, 2025
View reviewed changes
Copy link
Contributor

@AlexVlx AlexVlx left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM, thanks.

@arsenm Graphite App
Copy link
Contributor Author

arsenm commented Mar 6, 2025 •
edited
Loading

Merge activity

  • Mar 6, 9:39 AM EST: A user started a stack merge that includes this pull request via Graphite.
  • Mar 6, 9:40 AM EST: A user merged this pull request with Graphite.

@arsenm arsenm merged commit 4703f8b into main Mar 6, 2025
17 checks passed
@arsenm arsenm deleted the users/arsenm/clang/hip-avoid-ocml-log-exp branch March 6, 2025 14:40
This was referenced Nov 17, 2025
Closed
Closed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@AlexVlx AlexVlx AlexVlx approved these changes

@yxsamliu yxsamliu Awaiting requested review from yxsamliu

Assignees

No one assigned

Labels

backend:X86 clang:headers Headers provided by Clang, e.g. for intrinsics clang Clang issues not falling into any other category

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@arsenm @llvmbot @AlexVlx