Skip to content

Commit 160ae96

Browse files
author
Salinas, David
authored
Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V … (llvm#2273)
2 parents efd85c7 + 37411f5 commit 160ae96

File tree

3 files changed

+1680
-36
lines changed

3 files changed

+1680
-36
lines changed

‎clang/lib/Headers/__clang_hip_libdevice_declares.h‎

Lines changed: 12 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,8 @@
1414
#include "hip/hip_version.h"
1515
#endif // __has_include("hip/hip_version.h")
1616

17+
#define __PRIVATE_AS __attribute__((opencl_private))
18+
1719
#ifdef __cplusplus
1820
extern "C" {
1921
#endif
@@ -55,8 +57,7 @@ __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
5557
__device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
5658
__device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
5759
float);
58-
__device__ float __ocml_frexp_f32(float,
59-
__attribute__((address_space(5))) int *);
60+
__device__ float __ocml_frexp_f32(float, __PRIVATE_AS int *);
6061
__device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
6162
__device__ __attribute__((const)) int __ocml_ilogb_f32(float);
6263
__device__ __attribute__((const)) int __ocml_isfinite_f32(float);
@@ -74,8 +75,7 @@ __device__ __attribute__((pure)) float __ocml_native_log2_f32(float);
7475
__device__ __attribute__((const)) float __ocml_logb_f32(float);
7576
__device__ __attribute__((pure)) float __ocml_log_f32(float);
7677
__device__ __attribute__((pure)) float __ocml_native_log_f32(float);
77-
__device__ float __ocml_modf_f32(float,
78-
__attribute__((address_space(5))) float *);
78+
__device__ float __ocml_modf_f32(float, __PRIVATE_AS float *);
7979
__device__ __attribute__((const)) float __ocml_nearbyint_f32(float);
8080
__device__ __attribute__((const)) float __ocml_nextafter_f32(float, float);
8181
__device__ __attribute__((const)) float __ocml_len3_f32(float, float, float);
@@ -87,8 +87,7 @@ __device__ __attribute__((pure)) float __ocml_pow_f32(float, float);
8787
__device__ __attribute__((pure)) float __ocml_pown_f32(float, int);
8888
__device__ __attribute__((pure)) float __ocml_rcbrt_f32(float);
8989
__device__ __attribute__((const)) float __ocml_remainder_f32(float, float);
90-
__device__ float __ocml_remquo_f32(float, float,
91-
__attribute__((address_space(5))) int *);
90+
__device__ float __ocml_remquo_f32(float, float, __PRIVATE_AS int *);
9291
__device__ __attribute__((const)) float __ocml_rhypot_f32(float, float);
9392
__device__ __attribute__((const)) float __ocml_rint_f32(float);
9493
__device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float);
@@ -99,10 +98,8 @@ __device__ __attribute__((pure)) float __ocml_rsqrt_f32(float);
9998
__device__ __attribute__((const)) float __ocml_scalb_f32(float, float);
10099
__device__ __attribute__((const)) float __ocml_scalbn_f32(float, int);
101100
__device__ __attribute__((const)) int __ocml_signbit_f32(float);
102-
__device__ float __ocml_sincos_f32(float,
103-
__attribute__((address_space(5))) float *);
104-
__device__ float __ocml_sincospi_f32(float,
105-
__attribute__((address_space(5))) float *);
101+
__device__ float __ocml_sincos_f32(float, __PRIVATE_AS float *);
102+
__device__ float __ocml_sincospi_f32(float, __PRIVATE_AS float *);
106103
__device__ float __ocml_sin_f32(float);
107104
__device__ float __ocml_native_sin_f32(float);
108105
__device__ __attribute__((pure)) float __ocml_sinh_f32(float);
@@ -176,8 +173,7 @@ __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
176173
__device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
177174
__device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
178175
__device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
179-
__device__ double __ocml_frexp_f64(double,
180-
__attribute__((address_space(5))) int *);
176+
__device__ double __ocml_frexp_f64(double, __PRIVATE_AS int *);
181177
__device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
182178
__device__ __attribute__((const)) int __ocml_ilogb_f64(double);
183179
__device__ __attribute__((const)) int __ocml_isfinite_f64(double);
@@ -192,8 +188,7 @@ __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
192188
__device__ __attribute__((pure)) double __ocml_log2_f64(double);
193189
__device__ __attribute__((const)) double __ocml_logb_f64(double);
194190
__device__ __attribute__((pure)) double __ocml_log_f64(double);
195-
__device__ double __ocml_modf_f64(double,
196-
__attribute__((address_space(5))) double *);
191+
__device__ double __ocml_modf_f64(double, __PRIVATE_AS double *);
197192
__device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
198193
__device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
199194
__device__ __attribute__((const)) double __ocml_len3_f64(double, double,
@@ -206,8 +201,7 @@ __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
206201
__device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
207202
__device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
208203
__device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
209-
__device__ double __ocml_remquo_f64(double, double,
210-
__attribute__((address_space(5))) int *);
204+
__device__ double __ocml_remquo_f64(double, double, __PRIVATE_AS int *);
211205
__device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
212206
__device__ __attribute__((const)) double __ocml_rint_f64(double);
213207
__device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
@@ -219,10 +213,8 @@ __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
219213
__device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
220214
__device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
221215
__device__ __attribute__((const)) int __ocml_signbit_f64(double);
222-
__device__ double __ocml_sincos_f64(double,
223-
__attribute__((address_space(5))) double *);
224-
__device__ double
225-
__ocml_sincospi_f64(double, __attribute__((address_space(5))) double *);
216+
__device__ double __ocml_sincos_f64(double, __PRIVATE_AS double *);
217+
__device__ double __ocml_sincospi_f64(double, __PRIVATE_AS double *);
226218
__device__ double __ocml_sin_f64(double);
227219
__device__ __attribute__((pure)) double __ocml_sinh_f64(double);
228220
__device__ double __ocml_sinpi_f64(double);

‎clang/lib/Headers/__clang_hip_math.h‎

Lines changed: 12 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -55,6 +55,9 @@
5555
#define __DEVICE_NOCE__ __DEVICE__
5656
#endif
5757

58+
#pragma push_macro("__PRIVATE_AS")
59+
60+
#define __PRIVATE_AS __attribute__((opencl_private))
5861
// Device library provides fast low precision and slow full-recision
5962
// implementations for some functions. Which one gets selected depends on
6063
// __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
@@ -539,8 +542,7 @@ float modff(float __x, float *__iptr) {
539542
#ifdef __OPENMP_AMDGCN__
540543
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
541544
#endif
542-
float __r =
543-
__ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
545+
float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
544546
*__iptr = __tmp;
545547
return __r;
546548
}
@@ -625,8 +627,7 @@ float remquof(float __x, float __y, int *__quo) {
625627
#ifdef __OPENMP_AMDGCN__
626628
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
627629
#endif
628-
float __r = __ocml_remquo_f32(
629-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
630+
float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
630631
*__quo = __tmp;
631632

632633
return __r;
@@ -687,8 +688,7 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
687688
#ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
688689
__sincosf(__x, __sinptr, __cosptr);
689690
#else
690-
*__sinptr =
691-
__ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
691+
*__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
692692
*__cosptr = __tmp;
693693
#endif
694694
}
@@ -699,8 +699,7 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
699699
#ifdef __OPENMP_AMDGCN__
700700
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
701701
#endif
702-
*__sinptr = __ocml_sincospi_f32(
703-
__x, (__attribute__((address_space(5))) float *)&__tmp);
702+
*__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
704703
*__cosptr = __tmp;
705704
}
706705

@@ -943,8 +942,7 @@ double modf(double __x, double *__iptr) {
943942
#ifdef __OPENMP_AMDGCN__
944943
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
945944
#endif
946-
double __r =
947-
__ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
945+
double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
948946
*__iptr = __tmp;
949947

950948
return __r;
@@ -1037,8 +1035,7 @@ double remquo(double __x, double __y, int *__quo) {
10371035
#ifdef __OPENMP_AMDGCN__
10381036
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10391037
#endif
1040-
double __r = __ocml_remquo_f64(
1041-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
1038+
double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
10421039
*__quo = __tmp;
10431040

10441041
return __r;
@@ -1098,8 +1095,7 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
10981095
#ifdef __OPENMP_AMDGCN__
10991096
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
11001097
#endif
1101-
*__sinptr = __ocml_sincos_f64(
1102-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1098+
*__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
11031099
*__cosptr = __tmp;
11041100
}
11051101

@@ -1109,8 +1105,7 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
11091105
#ifdef __OPENMP_AMDGCN__
11101106
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
11111107
#endif
1112-
*__sinptr = __ocml_sincospi_f64(
1113-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1108+
*__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
11141109
*__cosptr = __tmp;
11151110
}
11161111

@@ -1358,6 +1353,7 @@ __host__ inline static int max(int __arg1, int __arg2) {
13581353

13591354
#pragma pop_macro("__DEVICE_NOCE__")
13601355
#pragma pop_macro("__DEVICE__")
1356+
#pragma pop_macro("__PRIVATE_AS")
13611357
#pragma pop_macro("__RETURN_TYPE")
13621358
#pragma pop_macro("__FAST_OR_SLOW")
13631359

0 commit comments

Comments
 (0)