Skip to content

Commit

Permalink
HIP: Directly call copysign builtins
Browse files Browse the repository at this point in the history
  • Loading branch information
arsenm committed Jun 8, 2023
1 parent c317a88 commit 4524a31
Show file tree
Hide file tree
Showing 2 changed files with 10 additions and 10 deletions.
4 changes: 2 additions & 2 deletions clang/lib/Headers/__clang_hip_math.h
Original file line number Diff line number Diff line change
Expand Up @@ -185,7 +185,7 @@ __DEVICE__
float ceilf(float __x) { return __ocml_ceil_f32(__x); }

__DEVICE__
float copysignf(float __x, float __y) { return __ocml_copysign_f32(__x, __y); }
float copysignf(float __x, float __y) { return __builtin_copysignf(__x, __y); }

__DEVICE__
float cosf(float __x) { return __ocml_cos_f32(__x); }
Expand Down Expand Up @@ -736,7 +736,7 @@ double ceil(double __x) { return __ocml_ceil_f64(__x); }

__DEVICE__
double copysign(double __x, double __y) {
return __ocml_copysign_f64(__x, __y);
return __builtin_copysign(__x, __y);
}

__DEVICE__
Expand Down
16 changes: 8 additions & 8 deletions clang/test/Headers/__clang_hip_math.hip
Original file line number Diff line number Diff line change
Expand Up @@ -503,27 +503,27 @@ extern "C" __device__ double test_ceil(double x) {

// DEFAULT-LABEL: @test_copysignf(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_copysign_f32(float noundef [[X:%.*]], float noundef [[Y:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: ret float [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: @test_copysignf(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_copysign_f32(float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret float [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.copysign.f32(float [[X:%.*]], float [[Y:%.*]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_copysignf(float x, float y) {
return copysignf(x, y);
}

// DEFAULT-LABEL: @test_copysign(
// DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_copysign_f64(double noundef [[X:%.*]], double noundef [[Y:%.*]]) #[[ATTR14]]
// DEFAULT-NEXT: ret double [[CALL_I]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: @test_copysign(
// FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_copysign_f64(double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) #[[ATTR14]]
// FINITEONLY-NEXT: ret double [[CALL_I]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.copysign.f64(double [[X:%.*]], double [[Y:%.*]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_copysign(double x, double y) {
return copysign(x, y);
Expand Down

0 comments on commit 4524a31

Please sign in to comment.