Skip to content

Commit 184e0f3

Browse files
AlexVlxAlexVlx
committed
"Generalise" math function signatures taking an explicit AS.
1 parent 91c0aa5 commit 184e0f3

File tree

3 files changed

+1687
-36
lines changed

3 files changed

+1687
-36
lines changed

‎clang/lib/Headers/__clang_hip_libdevice_declares.h‎

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

17+
#ifdef __SPIRV__
18+
#define __PRIVATE_AS __attribute__((address_space(0)))
19+
#else
20+
#define __PRIVATE_AS __attribute__((address_space(5)))
21+
#endif
22+
1723
#ifdef __cplusplus
1824
extern "C" {
1925
#endif
@@ -55,8 +61,7 @@ __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
5561
__device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
5662
__device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
5763
float);
58-
__device__ float __ocml_frexp_f32(float,
59-
__attribute__((address_space(5))) int *);
64+
__device__ float __ocml_frexp_f32(float, __PRIVATE_AS int *);
6065
__device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
6166
__device__ __attribute__((const)) int __ocml_ilogb_f32(float);
6267
__device__ __attribute__((const)) int __ocml_isfinite_f32(float);
@@ -74,8 +79,7 @@ __device__ __attribute__((pure)) float __ocml_native_log2_f32(float);
7479
__device__ __attribute__((const)) float __ocml_logb_f32(float);
7580
__device__ __attribute__((pure)) float __ocml_log_f32(float);
7681
__device__ __attribute__((pure)) float __ocml_native_log_f32(float);
77-
__device__ float __ocml_modf_f32(float,
78-
__attribute__((address_space(5))) float *);
82+
__device__ float __ocml_modf_f32(float, __PRIVATE_AS float *);
7983
__device__ __attribute__((const)) float __ocml_nearbyint_f32(float);
8084
__device__ __attribute__((const)) float __ocml_nextafter_f32(float, float);
8185
__device__ __attribute__((const)) float __ocml_len3_f32(float, float, float);
@@ -87,8 +91,7 @@ __device__ __attribute__((pure)) float __ocml_pow_f32(float, float);
8791
__device__ __attribute__((pure)) float __ocml_pown_f32(float, int);
8892
__device__ __attribute__((pure)) float __ocml_rcbrt_f32(float);
8993
__device__ __attribute__((const)) float __ocml_remainder_f32(float, float);
90-
__device__ float __ocml_remquo_f32(float, float,
91-
__attribute__((address_space(5))) int *);
94+
__device__ float __ocml_remquo_f32(float, float, __PRIVATE_AS int *);
9295
__device__ __attribute__((const)) float __ocml_rhypot_f32(float, float);
9396
__device__ __attribute__((const)) float __ocml_rint_f32(float);
9497
__device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float);
@@ -99,10 +102,8 @@ __device__ __attribute__((pure)) float __ocml_rsqrt_f32(float);
99102
__device__ __attribute__((const)) float __ocml_scalb_f32(float, float);
100103
__device__ __attribute__((const)) float __ocml_scalbn_f32(float, int);
101104
__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 *);
105+
__device__ float __ocml_sincos_f32(float, __PRIVATE_AS float *);
106+
__device__ float __ocml_sincospi_f32(float, __PRIVATE_AS float *);
106107
__device__ float __ocml_sin_f32(float);
107108
__device__ float __ocml_native_sin_f32(float);
108109
__device__ __attribute__((pure)) float __ocml_sinh_f32(float);
@@ -176,8 +177,7 @@ __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
176177
__device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
177178
__device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
178179
__device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
179-
__device__ double __ocml_frexp_f64(double,
180-
__attribute__((address_space(5))) int *);
180+
__device__ double __ocml_frexp_f64(double, __PRIVATE_AS int *);
181181
__device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
182182
__device__ __attribute__((const)) int __ocml_ilogb_f64(double);
183183
__device__ __attribute__((const)) int __ocml_isfinite_f64(double);
@@ -192,8 +192,7 @@ __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
192192
__device__ __attribute__((pure)) double __ocml_log2_f64(double);
193193
__device__ __attribute__((const)) double __ocml_logb_f64(double);
194194
__device__ __attribute__((pure)) double __ocml_log_f64(double);
195-
__device__ double __ocml_modf_f64(double,
196-
__attribute__((address_space(5))) double *);
195+
__device__ double __ocml_modf_f64(double, __PRIVATE_AS double *);
197196
__device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
198197
__device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
199198
__device__ __attribute__((const)) double __ocml_len3_f64(double, double,
@@ -206,8 +205,7 @@ __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
206205
__device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
207206
__device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
208207
__device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
209-
__device__ double __ocml_remquo_f64(double, double,
210-
__attribute__((address_space(5))) int *);
208+
__device__ double __ocml_remquo_f64(double, double, __PRIVATE_AS int *);
211209
__device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
212210
__device__ __attribute__((const)) double __ocml_rint_f64(double);
213211
__device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
@@ -219,10 +217,8 @@ __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
219217
__device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
220218
__device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
221219
__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 *);
220+
__device__ double __ocml_sincos_f64(double, __PRIVATE_AS double *);
221+
__device__ double __ocml_sincospi_f64(double, __PRIVATE_AS double *);
226222
__device__ double __ocml_sin_f64(double);
227223
__device__ __attribute__((pure)) double __ocml_sinh_f64(double);
228224
__device__ double __ocml_sinpi_f64(double);

‎clang/lib/Headers/__clang_hip_math.h‎

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

36+
#pragma push_macro("__PRIVATE_AS")
37+
38+
#ifdef __SPIRV__
39+
#define __PRIVATE_AS __attribute__((address_space(0)))
40+
#else
41+
#define __PRIVATE_AS __attribute__((address_space(5)))
42+
#endif
3643
// Device library provides fast low precision and slow full-recision
3744
// implementations for some functions. Which one gets selected depends on
3845
// __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
@@ -512,8 +519,7 @@ float modff(float __x, float *__iptr) {
512519
#ifdef __OPENMP_AMDGCN__
513520
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
514521
#endif
515-
float __r =
516-
__ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
522+
float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
517523
*__iptr = __tmp;
518524
return __r;
519525
}
@@ -595,8 +601,7 @@ float remquof(float __x, float __y, int *__quo) {
595601
#ifdef __OPENMP_AMDGCN__
596602
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
597603
#endif
598-
float __r = __ocml_remquo_f32(
599-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
604+
float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
600605
*__quo = __tmp;
601606

602607
return __r;
@@ -657,8 +662,7 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
657662
#ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
658663
__sincosf(__x, __sinptr, __cosptr);
659664
#else
660-
*__sinptr =
661-
__ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
665+
*__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
662666
*__cosptr = __tmp;
663667
#endif
664668
}
@@ -669,8 +673,7 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
669673
#ifdef __OPENMP_AMDGCN__
670674
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
671675
#endif
672-
*__sinptr = __ocml_sincospi_f32(
673-
__x, (__attribute__((address_space(5))) float *)&__tmp);
676+
*__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
674677
*__cosptr = __tmp;
675678
}
676679

@@ -913,8 +916,7 @@ double modf(double __x, double *__iptr) {
913916
#ifdef __OPENMP_AMDGCN__
914917
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
915918
#endif
916-
double __r =
917-
__ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
919+
double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
918920
*__iptr = __tmp;
919921

920922
return __r;
@@ -1004,8 +1006,7 @@ double remquo(double __x, double __y, int *__quo) {
10041006
#ifdef __OPENMP_AMDGCN__
10051007
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10061008
#endif
1007-
double __r = __ocml_remquo_f64(
1008-
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
1009+
double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
10091010
*__quo = __tmp;
10101011

10111012
return __r;
@@ -1065,8 +1066,7 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
10651066
#ifdef __OPENMP_AMDGCN__
10661067
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10671068
#endif
1068-
*__sinptr = __ocml_sincos_f64(
1069-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1069+
*__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
10701070
*__cosptr = __tmp;
10711071
}
10721072

@@ -1076,8 +1076,7 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
10761076
#ifdef __OPENMP_AMDGCN__
10771077
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
10781078
#endif
1079-
*__sinptr = __ocml_sincospi_f64(
1080-
__x, (__attribute__((address_space(5))) double *)&__tmp);
1079+
*__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
10811080
*__cosptr = __tmp;
10821081
}
10831082

@@ -1322,6 +1321,7 @@ __host__ inline static int max(int __arg1, int __arg2) {
13221321
#endif
13231322

13241323
#pragma pop_macro("__DEVICE__")
1324+
#pragma pop_macro("__PRIVATE_AS")
13251325
#pragma pop_macro("__RETURN_TYPE")
13261326
#pragma pop_macro("__FAST_OR_SLOW")
13271327

0 commit comments

Comments
 (0)