Skip to content

Commit 029bece

Browse files
AlexVlxAlexVlx
authored
[clang][HIP] Make some math not not work with AMDGCN SPIR-V (#128360)
Do not hardcode `address_space(5)` (`private`) in the ROCDL interface, as that breaks SPIRV generation (the latter uses 0). Add test. In the long run we should stop using ROCDL inline.
1 parent 9b6d0d7 commit 029bece

File tree

3 files changed

+1679
-36
lines changed

3 files changed

+1679
-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
@@ -33,6 +33,9 @@
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))
3639
// Device library provides fast low precision and slow full-recision
3740
// implementations for some functions. Which one gets selected depends on
3841
// __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
@@ -512,8 +515,7 @@ float modff(float __x, float *__iptr) {
512515
#ifdef __OPENMP_AMDGCN__
513516
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
514517
#endif
515-
float __r =
516-
__ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
518+
float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
517519
*__iptr = __tmp;
518520
return __r;
519521
}
@@ -595,8 +597,7 @@ float remquof(float __x, float __y, int *__quo) {
595597
#ifdef __OPENMP_AMDGCN__
596598
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
597599
#endif
598-
float __r = __ocml_remquo_f32(
599-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
600+
float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
600601
*__quo = __tmp;
601602

602603
return __r;
@@ -657,8 +658,7 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
657658
#ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
658659
__sincosf(__x, __sinptr, __cosptr);
659660
#else
660-
*__sinptr =
661-
__ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
661+
*__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
662662
*__cosptr = __tmp;
663663
#endif
664664
}
@@ -669,8 +669,7 @@ 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(
673-
__x, (__attribute__((address_space(5))) float *)&__tmp);
672+
*__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
674673
*__cosptr = __tmp;
675674
}
676675

@@ -913,8 +912,7 @@ double modf(double __x, double *__iptr) {
913912
#ifdef __OPENMP_AMDGCN__
914913
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
915914
#endif
916-
double __r =
917-
__ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
915+
double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
918916
*__iptr = __tmp;
919917

920918
return __r;
@@ -1004,8 +1002,7 @@ double remquo(double __x, double __y, int *__quo) {
10041002
#ifdef __OPENMP_AMDGCN__
10051003
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10061004
#endif
1007-
double __r = __ocml_remquo_f64(
1008-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
1005+
double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
10091006
*__quo = __tmp;
10101007

10111008
return __r;
@@ -1065,8 +1062,7 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
10651062
#ifdef __OPENMP_AMDGCN__
10661063
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10671064
#endif
1068-
*__sinptr = __ocml_sincos_f64(
1069-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1065+
*__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
10701066
*__cosptr = __tmp;
10711067
}
10721068

@@ -1076,8 +1072,7 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
10761072
#ifdef __OPENMP_AMDGCN__
10771073
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10781074
#endif
1079-
*__sinptr = __ocml_sincospi_f64(
1080-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1075+
*__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
10811076
*__cosptr = __tmp;
10821077
}
10831078

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

13241319
#pragma pop_macro("__DEVICE__")
1320+
#pragma pop_macro("__PRIVATE_AS")
13251321
#pragma pop_macro("__RETURN_TYPE")
13261322
#pragma pop_macro("__FAST_OR_SLOW")
13271323

0 commit comments

Comments
 (0)