Author: Matt Arsenault Date: 2023-06-28T15:06:54-04:00 New Revision: bf8e92c0e792cbe3c9cc50607a1e33c6912ffd0e
URL: https://github.com/llvm/llvm-project/commit/bf8e92c0e792cbe3c9cc50607a1e33c6912ffd0e DIFF: https://github.com/llvm/llvm-project/commit/bf8e92c0e792cbe3c9cc50607a1e33c6912ffd0e.diff LOG: HIP: Use frexp builtins in math headers Added: Modified: clang/lib/Headers/__clang_hip_math.h clang/test/Headers/__clang_hip_math.hip Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h index 7e95b66232122..26c2f77c82c63 100644 --- a/clang/lib/Headers/__clang_hip_math.h +++ b/clang/lib/Headers/__clang_hip_math.h @@ -257,8 +257,7 @@ float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); } __DEVICE__ float frexpf(float __x, int *__nptr) { - *__nptr = __builtin_amdgcn_frexp_expf(__x); - return __builtin_amdgcn_frexp_mantf(__x); + return __builtin_frexpf(__x, __nptr); } __DEVICE__ @@ -806,8 +805,7 @@ double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); } __DEVICE__ double frexp(double __x, int *__nptr) { - *__nptr = __builtin_amdgcn_frexp_exp(__x); - return __builtin_amdgcn_frexp_mant(__x); + return __builtin_frexp(__x, __nptr); } __DEVICE__ diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 984adf6da4ba2..e507aa3ccf4ef 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -1061,37 +1061,25 @@ extern "C" __device__ double test_fmod(double x, double y) { return fmod(x, y); } -// DEFAULT-LABEL: @test_frexpf( -// DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f32(float [[X:%.*]]) -// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12:![0-9]+]] -// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract float @llvm.amdgcn.frexp.mant.f32(float [[X]]) -// DEFAULT-NEXT: ret float [[TMP1]] -// -// FINITEONLY-LABEL: @test_frexpf( -// FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f32(float [[X:%.*]]) -// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12:![0-9]+]] -// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract float @llvm.amdgcn.frexp.mant.f32(float [[X]]) -// FINITEONLY-NEXT: ret float [[TMP1]] +// CHECK-LABEL: @test_frexpf( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[X:%.*]]) +// CHECK-NEXT: [[TMP1:%.*]] = extractvalue { float, i32 } [[TMP0]], 1 +// CHECK-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12:![0-9]+]] +// CHECK-NEXT: [[TMP2:%.*]] = extractvalue { float, i32 } [[TMP0]], 0 +// CHECK-NEXT: ret float [[TMP2]] // extern "C" __device__ float test_frexpf(float x, int* y) { return frexpf(x, y); } -// DEFAULT-LABEL: @test_frexp( -// DEFAULT-NEXT: entry: -// DEFAULT-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f64(double [[X:%.*]]) -// DEFAULT-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12]] -// DEFAULT-NEXT: [[TMP1:%.*]] = tail call contract double @llvm.amdgcn.frexp.mant.f64(double [[X]]) -// DEFAULT-NEXT: ret double [[TMP1]] -// -// FINITEONLY-LABEL: @test_frexp( -// FINITEONLY-NEXT: entry: -// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.frexp.exp.i32.f64(double [[X:%.*]]) -// FINITEONLY-NEXT: store i32 [[TMP0]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12]] -// FINITEONLY-NEXT: [[TMP1:%.*]] = tail call nnan ninf contract double @llvm.amdgcn.frexp.mant.f64(double [[X]]) -// FINITEONLY-NEXT: ret double [[TMP1]] +// CHECK-LABEL: @test_frexp( +// CHECK-NEXT: entry: +// CHECK-NEXT: [[TMP0:%.*]] = tail call { double, i32 } @llvm.frexp.f64.i32(double [[X:%.*]]) +// CHECK-NEXT: [[TMP1:%.*]] = extractvalue { double, i32 } [[TMP0]], 1 +// CHECK-NEXT: store i32 [[TMP1]], ptr [[Y:%.*]], align 4, !tbaa [[TBAA12]] +// CHECK-NEXT: [[TMP2:%.*]] = extractvalue { double, i32 } [[TMP0]], 0 +// CHECK-NEXT: ret double [[TMP2]] // extern "C" __device__ double test_frexp(double x, int* y) { return frexp(x, y); _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
