Skip to content

Commit 616f1c8

Browse files
AlexVlxsearlmc1
AlexVlx
authored and
searlmc1
committed
Reapply "[clang][HIP] Make some math not not work with AMDGCN SPIR-V llvm#128360" (llvm#129306)
This reapplies llvm#128360, the only change being that the modified tests also checks for the availability of the SPIRV target.
1 parent 8141028 commit 616f1c8

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
@@ -51,6 +51,9 @@
5151
#define __DEVICE_NOCE__ __DEVICE__
5252
#endif
5353

54+
#pragma push_macro("__PRIVATE_AS")
55+
56+
#define __PRIVATE_AS __attribute__((opencl_private))
5457
// Device library provides fast low precision and slow full-recision
5558
// implementations for some functions. Which one gets selected depends on
5659
// __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
@@ -535,8 +538,7 @@ float modff(float __x, float *__iptr) {
535538
#ifdef __OPENMP_AMDGCN__
536539
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
537540
#endif
538-
float __r =
539-
__ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
541+
float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
540542
*__iptr = __tmp;
541543
return __r;
542544
}
@@ -621,8 +623,7 @@ float remquof(float __x, float __y, int *__quo) {
621623
#ifdef __OPENMP_AMDGCN__
622624
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
623625
#endif
624-
float __r = __ocml_remquo_f32(
625-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
626+
float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
626627
*__quo = __tmp;
627628

628629
return __r;
@@ -683,8 +684,7 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
683684
#ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
684685
__sincosf(__x, __sinptr, __cosptr);
685686
#else
686-
*__sinptr =
687-
__ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
687+
*__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
688688
*__cosptr = __tmp;
689689
#endif
690690
}
@@ -695,8 +695,7 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
695695
#ifdef __OPENMP_AMDGCN__
696696
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
697697
#endif
698-
*__sinptr = __ocml_sincospi_f32(
699-
__x, (__attribute__((address_space(5))) float *)&__tmp);
698+
*__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
700699
*__cosptr = __tmp;
701700
}
702701

@@ -939,8 +938,7 @@ double modf(double __x, double *__iptr) {
939938
#ifdef __OPENMP_AMDGCN__
940939
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
941940
#endif
942-
double __r =
943-
__ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
941+
double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
944942
*__iptr = __tmp;
945943

946944
return __r;
@@ -1033,8 +1031,7 @@ double remquo(double __x, double __y, int *__quo) {
10331031
#ifdef __OPENMP_AMDGCN__
10341032
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10351033
#endif
1036-
double __r = __ocml_remquo_f64(
1037-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
1034+
double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
10381035
*__quo = __tmp;
10391036

10401037
return __r;
@@ -1094,8 +1091,7 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
10941091
#ifdef __OPENMP_AMDGCN__
10951092
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10961093
#endif
1097-
*__sinptr = __ocml_sincos_f64(
1098-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1094+
*__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
10991095
*__cosptr = __tmp;
11001096
}
11011097

@@ -1105,8 +1101,7 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
11051101
#ifdef __OPENMP_AMDGCN__
11061102
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
11071103
#endif
1108-
*__sinptr = __ocml_sincospi_f64(
1109-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1104+
*__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
11101105
*__cosptr = __tmp;
11111106
}
11121107

@@ -1354,6 +1349,7 @@ __host__ inline static int max(int __arg1, int __arg2) {
13541349

13551350
#pragma pop_macro("__DEVICE_NOCE__")
13561351
#pragma pop_macro("__DEVICE__")
1352+
#pragma pop_macro("__PRIVATE_AS")
13571353
#pragma pop_macro("__RETURN_TYPE")
13581354
#pragma pop_macro("__FAST_OR_SLOW")
13591355

0 commit comments

Comments
 (0)