This is an archive of the discontinued LLVM Phabricator instance.

Differential D138396

HIP: Directly call signbit builtins
ClosedPublic

Authored by arsenm on Nov 20 2022, 8:49 PM.

Details

Reviewers
JonChesterfield
b-sumner
yaxunl

Diff Detail

Event Timeline

arsenm created this revision.Nov 20 2022, 8:49 PM
Herald added a project: Restricted Project. · View Herald TranscriptNov 20 2022, 8:49 PM
arsenm requested review of this revision.Nov 20 2022, 8:49 PM
Herald added a subscriber: wdng. · View Herald TranscriptNov 20 2022, 8:49 PM
Harbormaster completed remote builds in B198717: Diff 476801.Nov 21 2022, 9:10 AM
arsenm added a comment.Jun 6 2023, 3:04 PM

ping

yaxunl accepted this revision.Jun 7 2023, 5:53 AM

LGTM. Thanks.

This revision is now accepted and ready to land.Jun 7 2023, 5:53 AM
arsenm closed this revision.Jun 8 2023, 12:23 PM

b51ae6d31f86de2e47b1912eb045ea37a0c5de23

Revision Contents

PathSize
clang/
lib/
Headers/
4 lines
test/
Headers/
15 lines

Diff 476801

clang/lib/Headers/__clang_hip_math.h

Show First 20 Lines • Show All 489 Lines • ▼ Show 20 Linesfloat scalblnf(float __x, long int __n) {
return (__n < INT_MAX) ? __ocml_scalbn_f32(__x, __n) return (__n < INT_MAX) ? __ocml_scalbn_f32(__x, __n)
: __ocml_scalb_f32(__x, __n); : __ocml_scalb_f32(__x, __n);
} }
__DEVICE__ __DEVICE__
float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); } float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __signbitf(float __x) { return __ocml_signbit_f32(__x); } __RETURN_TYPE __signbitf(float __x) { return __builtin_signbitf(__x); }
__DEVICE__ __DEVICE__
void sincosf(float __x, float *__sinptr, float *__cosptr) { void sincosf(float __x, float *__sinptr, float *__cosptr) {
float __tmp; float __tmp;
#ifdef __OPENMP_AMDGCN__ #ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc) #pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif #endif
*__sinptr = *__sinptr =
▲ Show 20 Lines • Show All 545 Lines • ▼ Show 20 Lines
double scalbln(double __x, long int __n) { double scalbln(double __x, long int __n) {
return (__n < INT_MAX) ? __ocml_scalbn_f64(__x, __n) return (__n < INT_MAX) ? __ocml_scalbn_f64(__x, __n)
: __ocml_scalb_f64(__x, __n); : __ocml_scalb_f64(__x, __n);
} }
__DEVICE__ __DEVICE__
double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); } double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); }
__DEVICE__ __DEVICE__
__RETURN_TYPE __signbit(double __x) { return __ocml_signbit_f64(__x); } __RETURN_TYPE __signbit(double __x) { return __builtin_signbit(__x); }
__DEVICE__ __DEVICE__
double sin(double __x) { return __ocml_sin_f64(__x); } double sin(double __x) { return __ocml_sin_f64(__x); }
__DEVICE__ __DEVICE__
void sincos(double __x, double *__sinptr, double *__cosptr) { void sincos(double __x, double *__sinptr, double *__cosptr) {
double __tmp; double __tmp;
#ifdef __OPENMP_AMDGCN__ #ifdef __OPENMP_AMDGCN__
▲ Show 20 Lines • Show All 262 Lines • Show Last 20 Lines

clang/test/Headers/__clang_hip_math.hip

Show First 20 Lines • Show All 2,718 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[CALL_I]]
// //
extern "C" __device__ double test_scalbn(double x, int y) { extern "C" __device__ double test_scalbn(double x, int y) {
return scalbn(x, y); return scalbn(x, y);
} }
// CHECK-LABEL: @test___signbitf( // CHECK-LABEL: @test___signbitf(
// CHECK-NEXT: entry: // CHECK-NEXT: entry:
// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f32(float noundef [[X:%.*]]) #[[ATTR15]] // CHECK-NEXT: [[TMP0:%.*]] = bitcast float [[X:%.*]] to i32
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // CHECK-NEXT: [[DOTLOBIT:%.*]] = lshr i32 [[TMP0]], 31
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // CHECK-NEXT: ret i32 [[DOTLOBIT]]
// CHECK-NEXT: ret i32 [[CONV]]
// //
extern "C" __device__ BOOL_TYPE test___signbitf(float x) { extern "C" __device__ BOOL_TYPE test___signbitf(float x) {
return __signbitf(x); return __signbitf(x);
} }
// CHECK-LABEL: @test___signbit( // CHECK-LABEL: @test___signbit(
// CHECK-NEXT: entry: // CHECK-NEXT: entry:
// CHECK-NEXT: [[CALL_I:%.*]] = tail call i32 @__ocml_signbit_f64(double noundef [[X:%.*]]) #[[ATTR15]] // CHECK-NEXT: [[TMP0:%.*]] = bitcast double [[X:%.*]] to i64
// CHECK-NEXT: [[TOBOOL_I:%.*]] = icmp ne i32 [[CALL_I]], 0 // CHECK-NEXT: [[DOTLOBIT:%.*]] = lshr i64 [[TMP0]], 63
// CHECK-NEXT: [[CONV:%.*]] = zext i1 [[TOBOOL_I]] to i32 // CHECK-NEXT: [[TMP1:%.*]] = trunc i64 [[DOTLOBIT]] to i32
// CHECK-NEXT: ret i32 [[CONV]] // CHECK-NEXT: ret i32 [[TMP1]]
// //
extern "C" __device__ BOOL_TYPE test___signbit(double x) { extern "C" __device__ BOOL_TYPE test___signbit(double x) {
return __signbit(x); return __signbit(x);
} }
// DEFAULT-LABEL: @test_sincosf( // DEFAULT-LABEL: @test_sincosf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5) // DEFAULT-NEXT: [[__TMP_I:%.*]] = alloca float, align 4, addrspace(5)
▲ Show 20 Lines • Show All 913 Lines • Show Last 20 Lines