Skip to content

[clang][HIP] Make some math not not work with AMDGCN SPIR-V #128360

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Feb 28, 2025
Merged

[clang][HIP] Make some math not not work with AMDGCN SPIR-V #128360

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
184e0f3
"Generalise" math function signatures taking an explicit AS.
AlexVlx Feb 22, 2025
8ade6ad
Merge branch 'main' of https://github.com/llvm/llvm-project into rocd…
AlexVlx Feb 28, 2025
b53179a
Apply suggestion.
AlexVlx Feb 28, 2025
4b012cd
Add missing paired attribute.
AlexVlx Feb 28, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
36 changes: 16 additions & 20 deletions clang/lib/Headers/__clang_hip_libdevice_declares.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,6 +14,12 @@
#include "hip/hip_version.h"
#endif // __has_include("hip/hip_version.h")

#ifdef __SPIRV__
#define __PRIVATE_AS __attribute__((address_space(0)))
Copy link
Contributor

@arsenm arsenm Feb 22, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Never use numbered address spaces. Unconditionally use __attribute__((opencl_private))

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Acknowledged, but the header was already using numbered address spaces (ROCDL is also using numbered rather than generic), and I really don't want to take on fixing the general badness living here/there. We shouldn't be using explicit ASes at all in what is supposed to be a generic interface, but that's a different kettle of fish. I'd like to keep this simple and land it, if possible, within the backport window as it's a hard break for us at the moment.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Using opencl_private is the simpler fix. You are calling code declared and defined in OpenCL, with the private address. Without a specific example of why this does not work out of the box, I insist you just go straight to using it

#else
#define __PRIVATE_AS __attribute__((address_space(5)))
#endif
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
#ifdef __SPIRV__
#define __PRIVATE_AS __attribute__((address_space(0)))
#else
#define __PRIVATE_AS __attribute__((address_space(5)))
#endif
#define __PRIVATE_AS __attribute__((opencl_private))


#ifdef __cplusplus
extern "C" {
#endif
Expand Down Expand Up @@ -55,8 +61,7 @@ __device__ __attribute__((const)) float __ocml_fmax_f32(float, float);
__device__ __attribute__((const)) float __ocml_fmin_f32(float, float);
__device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
float);
__device__ float __ocml_frexp_f32(float,
__attribute__((address_space(5))) int *);
__device__ float __ocml_frexp_f32(float, __PRIVATE_AS int *);
__device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
__device__ __attribute__((const)) int __ocml_ilogb_f32(float);
__device__ __attribute__((const)) int __ocml_isfinite_f32(float);
Expand All @@ -74,8 +79,7 @@ __device__ __attribute__((pure)) float __ocml_native_log2_f32(float);
__device__ __attribute__((const)) float __ocml_logb_f32(float);
__device__ __attribute__((pure)) float __ocml_log_f32(float);
__device__ __attribute__((pure)) float __ocml_native_log_f32(float);
__device__ float __ocml_modf_f32(float,
__attribute__((address_space(5))) float *);
__device__ float __ocml_modf_f32(float, __PRIVATE_AS float *);
__device__ __attribute__((const)) float __ocml_nearbyint_f32(float);
__device__ __attribute__((const)) float __ocml_nextafter_f32(float, float);
__device__ __attribute__((const)) float __ocml_len3_f32(float, float, float);
Expand All @@ -87,8 +91,7 @@ __device__ __attribute__((pure)) float __ocml_pow_f32(float, float);
__device__ __attribute__((pure)) float __ocml_pown_f32(float, int);
__device__ __attribute__((pure)) float __ocml_rcbrt_f32(float);
__device__ __attribute__((const)) float __ocml_remainder_f32(float, float);
__device__ float __ocml_remquo_f32(float, float,
__attribute__((address_space(5))) int *);
__device__ float __ocml_remquo_f32(float, float, __PRIVATE_AS int *);
__device__ __attribute__((const)) float __ocml_rhypot_f32(float, float);
__device__ __attribute__((const)) float __ocml_rint_f32(float);
__device__ __attribute__((const)) float __ocml_rlen3_f32(float, float, float);
Expand All @@ -99,10 +102,8 @@ __device__ __attribute__((pure)) float __ocml_rsqrt_f32(float);
__device__ __attribute__((const)) float __ocml_scalb_f32(float, float);
__device__ __attribute__((const)) float __ocml_scalbn_f32(float, int);
__device__ __attribute__((const)) int __ocml_signbit_f32(float);
__device__ float __ocml_sincos_f32(float,
__attribute__((address_space(5))) float *);
__device__ float __ocml_sincospi_f32(float,
__attribute__((address_space(5))) float *);
__device__ float __ocml_sincos_f32(float, __PRIVATE_AS float *);
__device__ float __ocml_sincospi_f32(float, __PRIVATE_AS float *);
__device__ float __ocml_sin_f32(float);
__device__ float __ocml_native_sin_f32(float);
__device__ __attribute__((pure)) float __ocml_sinh_f32(float);
Expand Down Expand Up @@ -176,8 +177,7 @@ __device__ __attribute__((const)) double __ocml_fma_f64(double, double, double);
__device__ __attribute__((const)) double __ocml_fmax_f64(double, double);
__device__ __attribute__((const)) double __ocml_fmin_f64(double, double);
__device__ __attribute__((const)) double __ocml_fmod_f64(double, double);
__device__ double __ocml_frexp_f64(double,
__attribute__((address_space(5))) int *);
__device__ double __ocml_frexp_f64(double, __PRIVATE_AS int *);
__device__ __attribute__((const)) double __ocml_hypot_f64(double, double);
__device__ __attribute__((const)) int __ocml_ilogb_f64(double);
__device__ __attribute__((const)) int __ocml_isfinite_f64(double);
Expand All @@ -192,8 +192,7 @@ __device__ __attribute__((pure)) double __ocml_log1p_f64(double);
__device__ __attribute__((pure)) double __ocml_log2_f64(double);
__device__ __attribute__((const)) double __ocml_logb_f64(double);
__device__ __attribute__((pure)) double __ocml_log_f64(double);
__device__ double __ocml_modf_f64(double,
__attribute__((address_space(5))) double *);
__device__ double __ocml_modf_f64(double, __PRIVATE_AS double *);
__device__ __attribute__((const)) double __ocml_nearbyint_f64(double);
__device__ __attribute__((const)) double __ocml_nextafter_f64(double, double);
__device__ __attribute__((const)) double __ocml_len3_f64(double, double,
Expand All @@ -206,8 +205,7 @@ __device__ __attribute__((pure)) double __ocml_pow_f64(double, double);
__device__ __attribute__((pure)) double __ocml_pown_f64(double, int);
__device__ __attribute__((pure)) double __ocml_rcbrt_f64(double);
__device__ __attribute__((const)) double __ocml_remainder_f64(double, double);
__device__ double __ocml_remquo_f64(double, double,
__attribute__((address_space(5))) int *);
__device__ double __ocml_remquo_f64(double, double, __PRIVATE_AS int *);
__device__ __attribute__((const)) double __ocml_rhypot_f64(double, double);
__device__ __attribute__((const)) double __ocml_rint_f64(double);
__device__ __attribute__((const)) double __ocml_rlen3_f64(double, double,
Expand All @@ -219,10 +217,8 @@ __device__ __attribute__((pure)) double __ocml_rsqrt_f64(double);
__device__ __attribute__((const)) double __ocml_scalb_f64(double, double);
__device__ __attribute__((const)) double __ocml_scalbn_f64(double, int);
__device__ __attribute__((const)) int __ocml_signbit_f64(double);
__device__ double __ocml_sincos_f64(double,
__attribute__((address_space(5))) double *);
__device__ double
__ocml_sincospi_f64(double, __attribute__((address_space(5))) double *);
__device__ double __ocml_sincos_f64(double, __PRIVATE_AS double *);
__device__ double __ocml_sincospi_f64(double, __PRIVATE_AS double *);
__device__ double __ocml_sin_f64(double);
__device__ __attribute__((pure)) double __ocml_sinh_f64(double);
__device__ double __ocml_sinpi_f64(double);
Expand Down
32 changes: 16 additions & 16 deletions clang/lib/Headers/__clang_hip_math.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,6 +33,13 @@
#define __DEVICE__ static __device__ inline __attribute__((always_inline))
#endif

#pragma push_macro("__PRIVATE_AS")

#ifdef __SPIRV__
#define __PRIVATE_AS __attribute__((address_space(0)))
#else
#define __PRIVATE_AS __attribute__((address_space(5)))
#endif
// Device library provides fast low precision and slow full-recision
// implementations for some functions. Which one gets selected depends on
// __CLANG_GPU_APPROX_TRANSCENDENTALS__ which gets defined by clang if
Expand Down Expand Up @@ -512,8 +519,7 @@ float modff(float __x, float *__iptr) {
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
float __r =
__ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
float __r = __ocml_modf_f32(__x, (__PRIVATE_AS float *)&__tmp);
*__iptr = __tmp;
return __r;
}
Expand Down Expand Up @@ -595,8 +601,7 @@ float remquof(float __x, float __y, int *__quo) {
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
float __r = __ocml_remquo_f32(
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
float __r = __ocml_remquo_f32(__x, __y, (__PRIVATE_AS int *)&__tmp);
*__quo = __tmp;

return __r;
Expand Down Expand Up @@ -657,8 +662,7 @@ void sincosf(float __x, float *__sinptr, float *__cosptr) {
#ifdef __CLANG_CUDA_APPROX_TRANSCENDENTALS__
__sincosf(__x, __sinptr, __cosptr);
#else
*__sinptr =
__ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
*__sinptr = __ocml_sincos_f32(__x, (__PRIVATE_AS float *)&__tmp);
*__cosptr = __tmp;
#endif
}
Expand All @@ -669,8 +673,7 @@ void sincospif(float __x, float *__sinptr, float *__cosptr) {
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
*__sinptr = __ocml_sincospi_f32(
__x, (__attribute__((address_space(5))) float *)&__tmp);
*__sinptr = __ocml_sincospi_f32(__x, (__PRIVATE_AS float *)&__tmp);
*__cosptr = __tmp;
}

Expand Down Expand Up @@ -913,8 +916,7 @@ double modf(double __x, double *__iptr) {
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
double __r =
__ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
double __r = __ocml_modf_f64(__x, (__PRIVATE_AS double *)&__tmp);
*__iptr = __tmp;

return __r;
Expand Down Expand Up @@ -1004,8 +1006,7 @@ double remquo(double __x, double __y, int *__quo) {
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
double __r = __ocml_remquo_f64(
__x, __y, (__attribute__((address_space(5))) int *)&__tmp);
double __r = __ocml_remquo_f64(__x, __y, (__PRIVATE_AS int *)&__tmp);
*__quo = __tmp;

return __r;
Expand Down Expand Up @@ -1065,8 +1066,7 @@ void sincos(double __x, double *__sinptr, double *__cosptr) {
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
*__sinptr = __ocml_sincos_f64(
__x, (__attribute__((address_space(5))) double *)&__tmp);
*__sinptr = __ocml_sincos_f64(__x, (__PRIVATE_AS double *)&__tmp);
*__cosptr = __tmp;
}

Expand All @@ -1076,8 +1076,7 @@ void sincospi(double __x, double *__sinptr, double *__cosptr) {
#ifdef __OPENMP_AMDGCN__
#pragma omp allocate(__tmp) allocator(omp_thread_mem_alloc)
#endif
*__sinptr = __ocml_sincospi_f64(
__x, (__attribute__((address_space(5))) double *)&__tmp);
*__sinptr = __ocml_sincospi_f64(__x, (__PRIVATE_AS double *)&__tmp);
*__cosptr = __tmp;
}

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

#pragma pop_macro("__DEVICE__")
#pragma pop_macro("__PRIVATE_AS")
#pragma pop_macro("__RETURN_TYPE")
#pragma pop_macro("__FAST_OR_SLOW")

Expand Down
Loading