This is an archive of the discontinued LLVM Phabricator instance.

Differential D154000

HIP: Directly call round builtins
ClosedPublic

Authored by arsenm on Jun 28 2023, 1:49 PM.

Details

Reviewers
JonChesterfield
yaxunl
jhuber6
Summary

Pretty sure these lround->round cases are just wrong

Diff Detail

Event Timeline

arsenm created this revision.Jun 28 2023, 1:49 PM
Herald added a project: Restricted Project. · View Herald TranscriptJun 28 2023, 1:49 PM
arsenm requested review of this revision.Jun 28 2023, 1:49 PM
Herald added a subscriber: wdng. · View Herald TranscriptJun 28 2023, 1:49 PM
jhuber6 accepted this revision.Jun 28 2023, 1:53 PM
This revision is now accepted and ready to land.Jun 28 2023, 1:53 PM
Harbormaster completed remote builds in B241883: Diff 535509.Jun 28 2023, 3:08 PM
arsenm closed this revision.Jul 7 2023, 6:29 AM

9a317516a515f3c5b15f9060329a503e8f261c7f

Revision Contents

PathSize
clang/
lib/
Headers/
12 lines
test/
Headers/
48 lines

Diff 535509

clang/lib/Headers/__clang_hip_math.h

Show First 20 Lines • Show All 306 Lines • ▼ Show 20 Lines
__DEVICE__ __DEVICE__
float lgammaf(float __x) { return __ocml_lgamma_f32(__x); } float lgammaf(float __x) { return __ocml_lgamma_f32(__x); }
__DEVICE__ __DEVICE__
long long int llrintf(float __x) { return __ocml_rint_f32(__x); } long long int llrintf(float __x) { return __ocml_rint_f32(__x); }
__DEVICE__ __DEVICE__
long long int llroundf(float __x) { return __ocml_round_f32(__x); } long long int llroundf(float __x) { return __builtin_roundf(__x); }
__DEVICE__ __DEVICE__
float log10f(float __x) { return __ocml_log10_f32(__x); } float log10f(float __x) { return __ocml_log10_f32(__x); }
__DEVICE__ __DEVICE__
float log1pf(float __x) { return __ocml_log1p_f32(__x); } float log1pf(float __x) { return __ocml_log1p_f32(__x); }
__DEVICE__ __DEVICE__
float log2f(float __x) { return __ocml_log2_f32(__x); } float log2f(float __x) { return __ocml_log2_f32(__x); }
__DEVICE__ __DEVICE__
float logbf(float __x) { return __ocml_logb_f32(__x); } float logbf(float __x) { return __ocml_logb_f32(__x); }
__DEVICE__ __DEVICE__
float logf(float __x) { return __ocml_log_f32(__x); } float logf(float __x) { return __ocml_log_f32(__x); }
__DEVICE__ __DEVICE__
long int lrintf(float __x) { return __ocml_rint_f32(__x); } long int lrintf(float __x) { return __ocml_rint_f32(__x); }
__DEVICE__ __DEVICE__
long int lroundf(float __x) { return __ocml_round_f32(__x); } long int lroundf(float __x) { return __builtin_roundf(__x); }
__DEVICE__ __DEVICE__
float modff(float __x, float *__iptr) { float modff(float __x, float *__iptr) {
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
float __r = float __r =
▲ Show 20 Lines • Show All 110 Lines • ▼ Show 20 Lines while (__dim--) {
__r += __a[0] * __a[0]; __r += __a[0] * __a[0];
++__a; ++__a;
} }
return __ocml_rsqrt_f32(__r); return __ocml_rsqrt_f32(__r);
} }
__DEVICE__ __DEVICE__
float roundf(float __x) { return __ocml_round_f32(__x); } float roundf(float __x) { return __builtin_roundf(__x); }
__DEVICE__ __DEVICE__
float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); } float rsqrtf(float __x) { return __ocml_rsqrt_f32(__x); }
__DEVICE__ __DEVICE__
float scalblnf(float __x, long int __n) { float scalblnf(float __x, long int __n) {
return (__n < INT_MAX) ? __builtin_amdgcn_ldexpf(__x, __n) return (__n < INT_MAX) ? __builtin_amdgcn_ldexpf(__x, __n)
: __ocml_scalb_f32(__x, __n); : __ocml_scalb_f32(__x, __n);
▲ Show 20 Lines • Show All 383 Lines • ▼ Show 20 Lines
__DEVICE__ __DEVICE__
double lgamma(double __x) { return __ocml_lgamma_f64(__x); } double lgamma(double __x) { return __ocml_lgamma_f64(__x); }
__DEVICE__ __DEVICE__
long long int llrint(double __x) { return __ocml_rint_f64(__x); } long long int llrint(double __x) { return __ocml_rint_f64(__x); }
__DEVICE__ __DEVICE__
long long int llround(double __x) { return __ocml_round_f64(__x); } long long int llround(double __x) { return __builtin_round(__x); }
__DEVICE__ __DEVICE__
double log(double __x) { return __ocml_log_f64(__x); } double log(double __x) { return __ocml_log_f64(__x); }
__DEVICE__ __DEVICE__
double log10(double __x) { return __ocml_log10_f64(__x); } double log10(double __x) { return __ocml_log10_f64(__x); }
__DEVICE__ __DEVICE__
double log1p(double __x) { return __ocml_log1p_f64(__x); } double log1p(double __x) { return __ocml_log1p_f64(__x); }
__DEVICE__ __DEVICE__
double log2(double __x) { return __ocml_log2_f64(__x); } double log2(double __x) { return __ocml_log2_f64(__x); }
__DEVICE__ __DEVICE__
double logb(double __x) { return __ocml_logb_f64(__x); } double logb(double __x) { return __ocml_logb_f64(__x); }
__DEVICE__ __DEVICE__
long int lrint(double __x) { return __ocml_rint_f64(__x); } long int lrint(double __x) { return __ocml_rint_f64(__x); }
__DEVICE__ __DEVICE__
long int lround(double __x) { return __ocml_round_f64(__x); } long int lround(double __x) { return __builtin_round(__x); }
__DEVICE__ __DEVICE__
double modf(double __x, double *__iptr) { double modf(double __x, double *__iptr) {
double __tmp; double __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
double __r = double __r =
▲ Show 20 Lines • Show All 118 Lines • ▼ Show 20 Lines
} }
__DEVICE__ __DEVICE__
double rnorm4d(double __x, double __y, double __z, double __w) { double rnorm4d(double __x, double __y, double __z, double __w) {
return __ocml_rlen4_f64(__x, __y, __z, __w); return __ocml_rlen4_f64(__x, __y, __z, __w);
} }
__DEVICE__ __DEVICE__
double round(double __x) { return __ocml_round_f64(__x); } double round(double __x) { return __builtin_round(__x); }
__DEVICE__ __DEVICE__
double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); } double rsqrt(double __x) { return __ocml_rsqrt_f64(__x); }
__DEVICE__ __DEVICE__
double scalbln(double __x, long int __n) { double scalbln(double __x, long int __n) {
return (__n < INT_MAX) ? __builtin_amdgcn_ldexp(__x, __n) return (__n < INT_MAX) ? __builtin_amdgcn_ldexp(__x, __n)
: __ocml_scalb_f64(__x, __n); : __ocml_scalb_f64(__x, __n);
▲ Show 20 Lines • Show All 276 Lines • Show Last 20 Lines

clang/test/Headers/__clang_hip_math.hip

Show First 20 Lines • Show All 1,509 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret i64 [[CONV_I]] // FINITEONLY-NEXT: ret i64 [[CONV_I]]
// //
extern "C" __device__ long long int test_llrint(double x) { extern "C" __device__ long long int test_llrint(double x) {
return llrint(x); return llrint(x);
} }
// DEFAULT-LABEL: @test_llroundf( // DEFAULT-LABEL: @test_llroundf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]] // DEFAULT-NEXT: ret i64 [[CONV_I]]
// //
// FINITEONLY-LABEL: @test_llroundf( // FINITEONLY-LABEL: @test_llroundf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]] // FINITEONLY-NEXT: ret i64 [[CONV_I]]
// //
extern "C" __device__ long long int test_llroundf(float x) { extern "C" __device__ long long int test_llroundf(float x) {
return llroundf(x); return llroundf(x);
} }
// DEFAULT-LABEL: @test_llround( // DEFAULT-LABEL: @test_llround(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]] // DEFAULT-NEXT: ret i64 [[CONV_I]]
// //
// FINITEONLY-LABEL: @test_llround( // FINITEONLY-LABEL: @test_llround(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]] // FINITEONLY-NEXT: ret i64 [[CONV_I]]
// //
extern "C" __device__ long long int test_llround(double x) { extern "C" __device__ long long int test_llround(double x) {
return llround(x); return llround(x);
} }
// DEFAULT-LABEL: @test_log10f( // DEFAULT-LABEL: @test_log10f(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
▲ Show 20 Lines • Show All 136 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret i64 [[CONV_I]] // FINITEONLY-NEXT: ret i64 [[CONV_I]]
// //
extern "C" __device__ long int test_lrint(double x) { extern "C" __device__ long int test_lrint(double x) {
return lrint(x); return lrint(x);
} }
// DEFAULT-LABEL: @test_lroundf( // DEFAULT-LABEL: @test_lroundf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]] // DEFAULT-NEXT: ret i64 [[CONV_I]]
// //
// FINITEONLY-LABEL: @test_lroundf( // FINITEONLY-LABEL: @test_lroundf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[CALL_I]] to i64 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi float [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]] // FINITEONLY-NEXT: ret i64 [[CONV_I]]
// //
extern "C" __device__ long int test_lroundf(float x) { extern "C" __device__ long int test_lroundf(float x) {
return lroundf(x); return lroundf(x);
} }
// DEFAULT-LABEL: @test_lround( // DEFAULT-LABEL: @test_lround(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
// DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // DEFAULT-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// DEFAULT-NEXT: ret i64 [[CONV_I]] // DEFAULT-NEXT: ret i64 [[CONV_I]]
// //
// FINITEONLY-LABEL: @test_lround( // FINITEONLY-LABEL: @test_lround(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]])
// FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[CALL_I]] to i64 // FINITEONLY-NEXT: [[CONV_I:%.*]] = fptosi double [[TMP0]] to i64
// FINITEONLY-NEXT: ret i64 [[CONV_I]] // FINITEONLY-NEXT: ret i64 [[CONV_I]]
// //
extern "C" __device__ long int test_lround(double x) { extern "C" __device__ long int test_lround(double x) {
return lround(x); return lround(x);
} }
// DEFAULT-LABEL: @test_modff( // DEFAULT-LABEL: @test_modff(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
▲ Show 20 Lines • Show All 885 Lines • ▼ Show 20 Lines
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[CALL_I]]
// //
extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w) { extern "C" __device__ double test_rnorm4d(double x, double y, double z, double w) {
return rnorm4d(x, y, z, w); return rnorm4d(x, y, z, w);
} }
// DEFAULT-LABEL: @test_roundf( // DEFAULT-LABEL: @test_roundf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_round_f32(float noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.round.f32(float [[X:%.*]])
// DEFAULT-NEXT: ret float [[CALL_I]] // DEFAULT-NEXT: ret float [[TMP0]]
// //
// FINITEONLY-LABEL: @test_roundf( // FINITEONLY-LABEL: @test_roundf(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) float @__ocml_round_f32(float noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract float @llvm.round.f32(float [[X:%.*]])
// FINITEONLY-NEXT: ret float [[CALL_I]] // FINITEONLY-NEXT: ret float [[TMP0]]
// //
extern "C" __device__ float test_roundf(float x) { extern "C" __device__ float test_roundf(float x) {
return roundf(x); return roundf(x);
} }
// DEFAULT-LABEL: @test_round( // DEFAULT-LABEL: @test_round(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract double @__ocml_round_f64(double noundef [[X:%.*]]) #[[ATTR14]] // DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.round.f64(double [[X:%.*]])
// DEFAULT-NEXT: ret double [[CALL_I]] // DEFAULT-NEXT: ret double [[TMP0]]
// //
// FINITEONLY-LABEL: @test_round( // FINITEONLY-LABEL: @test_round(
// FINITEONLY-NEXT: entry: // FINITEONLY-NEXT: entry:
// FINITEONLY-NEXT: [[CALL_I:%.*]] = tail call nnan ninf contract nofpclass(nan inf) double @__ocml_round_f64(double noundef nofpclass(nan inf) [[X:%.*]]) #[[ATTR14]] // FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract double @llvm.round.f64(double [[X:%.*]])
// FINITEONLY-NEXT: ret double [[CALL_I]] // FINITEONLY-NEXT: ret double [[TMP0]]
// //
extern "C" __device__ double test_round(double x) { extern "C" __device__ double test_round(double x) {
return round(x); return round(x);
} }
// DEFAULT-LABEL: @test_rsqrtf( // DEFAULT-LABEL: @test_rsqrtf(
// DEFAULT-NEXT: entry: // DEFAULT-NEXT: entry:
// DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]] // DEFAULT-NEXT: [[CALL_I:%.*]] = tail call contract float @__ocml_rsqrt_f32(float noundef [[X:%.*]]) #[[ATTR15]]
▲ Show 20 Lines • Show All 1,072 Lines • Show Last 20 Lines