Skip to content

Commit 926600a

Browse files
AlexVlxAlexVlx
authored
Revert "[clang][HIP] Make some math not not work with AMDGCN SPIR-V" (llvm#129280)
Reverts llvm#128360 pending resolution of odd test break.
1 parent bdace10 commit 926600a

File tree

3 files changed

+36
-1679
lines changed

3 files changed

+36
-1679
lines changed

‎clang/lib/Headers/__clang_hip_libdevice_declares.h‎

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

17-
#define __PRIVATE_AS __attribute__((opencl_private))
18-
1917
#ifdef __cplusplus
2018
extern "C" {
2119
#endif
@@ -57,7 +55,8 @@ __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
5755
__device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
5856
__device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
5957
float);
60-
__device__ float __ocml_frexp_f32(float, __PRIVATE_AS int *);
58+
__device__ float __ocml_frexp_f32(float,
59+
__attribute__((address_space(5))) int *);
6160
__device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
6261
__device__ __attribute__((const)) int __ocml_ilogb_f32(float);
6362
__device__ __attribute__((const)) int __ocml_isfinite_f32(float);
@@ -75,7 +74,8 @@ __device__ __attribute__((pure)) float __ocml_native_log2_f32(float);
7574
__device__ __attribute__((const)) float __ocml_logb_f32(float);
7675
__device__ __attribute__((pure)) float __ocml_log_f32(float);
7776
__device__ __attribute__((pure)) float __ocml_native_log_f32(float);
78-
__device__ float __ocml_modf_f32(float, __PRIVATE_AS float *);
77+
__device__ float __ocml_modf_f32(float,
78+
__attribute__((address_space(5))) 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,7 +87,8 @@ __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, __PRIVATE_AS int *);
90+
__device__ float __ocml_remquo_f32(float, float,
91+
__attribute__((address_space(5))) int *);
9192
__device__ __attribute__((const)) float __ocml_rhypot_f32(float, float);
9293
__device__ __attribute__((const)) float __ocml_rint_f32(float);
9394
__device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float);
@@ -98,8 +99,10 @@ __device__ __attribute__((pure)) float __ocml_rsqrt_f32(float);
9899
__device__ __attribute__((const)) float __ocml_scalb_f32(float, float);
99100
__device__ __attribute__((const)) float __ocml_scalbn_f32(float, int);
100101
__device__ __attribute__((const)) int __ocml_signbit_f32(float);
101-
__device__ float __ocml_sincos_f32(float, __PRIVATE_AS float *);
102-
__device__ float __ocml_sincospi_f32(float, __PRIVATE_AS 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 *);
103106
__device__ float __ocml_sin_f32(float);
104107
__device__ float __ocml_native_sin_f32(float);
105108
__device__ __attribute__((pure)) float __ocml_sinh_f32(float);
@@ -173,7 +176,8 @@ __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
173176
__device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
174177
__device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
175178
__device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
176-
__device__ double __ocml_frexp_f64(double, __PRIVATE_AS int *);
179+
__device__ double __ocml_frexp_f64(double,
180+
__attribute__((address_space(5))) int *);
177181
__device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
178182
__device__ __attribute__((const)) int __ocml_ilogb_f64(double);
179183
__device__ __attribute__((const)) int __ocml_isfinite_f64(double);
@@ -188,7 +192,8 @@ __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
188192
__device__ __attribute__((pure)) double __ocml_log2_f64(double);
189193
__device__ __attribute__((const)) double __ocml_logb_f64(double);
190194
__device__ __attribute__((pure)) double __ocml_log_f64(double);
191-
__device__ double __ocml_modf_f64(double, __PRIVATE_AS double *);
195+
__device__ double __ocml_modf_f64(double,
196+
__attribute__((address_space(5))) double *);
192197
__device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
193198
__device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
194199
__device__ __attribute__((const)) double __ocml_len3_f64(double, double,
@@ -201,7 +206,8 @@ __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
201206
__device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
202207
__device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
203208
__device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
204-
__device__ double __ocml_remquo_f64(double, double, __PRIVATE_AS int *);
209+
__device__ double __ocml_remquo_f64(double, double,
210+
__attribute__((address_space(5))) int *);
205211
__device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
206212
__device__ __attribute__((const)) double __ocml_rint_f64(double);
207213
__device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
@@ -213,8 +219,10 @@ __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
213219
__device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
214220
__device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
215221
__device__ __attribute__((const)) int __ocml_signbit_f64(double);
216-
__device__ double __ocml_sincos_f64(double, __PRIVATE_AS double *);
217-
__device__ double __ocml_sincospi_f64(double, __PRIVATE_AS 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 *);
218226
__device__ double __ocml_sin_f64(double);
219227
__device__ __attribute__((pure)) double __ocml_sinh_f64(double);
220228
__device__ double __ocml_sinpi_f64(double);

‎clang/lib/Headers/__clang_hip_math.h‎

Lines changed: 16 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -33,9 +33,6 @@
3333
#define __DEVICE__ static __device__ inline __attribute__((always_inline))
3434
#endif
3535

36-
#pragma push_macro("__PRIVATE_AS")
37-
38-
#define __PRIVATE_AS __attribute__((opencl_private))
3936
// Device library provides fast low precision and slow full-recision
4037
// implementations for some functions. Which one gets selected depends on
4138
// __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
@@ -515,7 +512,8 @@ float modff(float __x, float *__iptr) {
515512
#ifdef __OPENMP_AMDGCN__
516513
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
517514
#endif
518-
float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
515+
float __r =
516+
__ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
519517
*__iptr = __tmp;
520518
return __r;
521519
}
@@ -597,7 +595,8 @@ float remquof(float __x, float __y, int *__quo) {
597595
#ifdef __OPENMP_AMDGCN__
598596
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
599597
#endif
600-
float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
598+
float __r = __ocml_remquo_f32(
599+
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
601600
*__quo = __tmp;
602601

603602
return __r;
@@ -658,7 +657,8 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
658657
#ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
659658
__sincosf(__x, __sinptr, __cosptr);
660659
#else
661-
*__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
660+
*__sinptr =
661+
__ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
662662
*__cosptr = __tmp;
663663
#endif
664664
}
@@ -669,7 +669,8 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
669669
#ifdef __OPENMP_AMDGCN__
670670
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
671671
#endif
672-
*__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
672+
*__sinptr = __ocml_sincospi_f32(
673+
__x, (__attribute__((address_space(5))) float *)&__tmp);
673674
*__cosptr = __tmp;
674675
}
675676

@@ -912,7 +913,8 @@ double modf(double __x, double *__iptr) {
912913
#ifdef __OPENMP_AMDGCN__
913914
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
914915
#endif
915-
double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
916+
double __r =
917+
__ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
916918
*__iptr = __tmp;
917919

918920
return __r;
@@ -1002,7 +1004,8 @@ double remquo(double __x, double __y, int *__quo) {
10021004
#ifdef __OPENMP_AMDGCN__
10031005
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10041006
#endif
1005-
double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
1007+
double __r = __ocml_remquo_f64(
1008+
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
10061009
*__quo = __tmp;
10071010

10081011
return __r;
@@ -1062,7 +1065,8 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
10621065
#ifdef __OPENMP_AMDGCN__
10631066
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10641067
#endif
1065-
*__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
1068+
*__sinptr = __ocml_sincos_f64(
1069+
__x, (__attribute__((address_space(5))) double *)&__tmp);
10661070
*__cosptr = __tmp;
10671071
}
10681072

@@ -1072,7 +1076,8 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
10721076
#ifdef __OPENMP_AMDGCN__
10731077
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10741078
#endif
1075-
*__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
1079+
*__sinptr = __ocml_sincospi_f64(
1080+
__x, (__attribute__((address_space(5))) double *)&__tmp);
10761081
*__cosptr = __tmp;
10771082
}
10781083

@@ -1317,7 +1322,6 @@ __host__ inline static int max(int __arg1, int __arg2) {
13171322
#endif
13181323

13191324
#pragma pop_macro("__DEVICE__")
1320-
#pragma pop_macro("__PRIVATE_AS")
13211325
#pragma pop_macro("__RETURN_TYPE")
13221326
#pragma pop_macro("__FAST_OR_SLOW")
13231327

0 commit comments

Comments
 (0)