diff --git a/libdevice/atomic.hpp b/libdevice/atomic.hpp index 3b6d1cf71f441..ca35fa8767cd0 100644 --- a/libdevice/atomic.hpp +++ b/libdevice/atomic.hpp @@ -11,7 +11,7 @@ #include "device.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #define SPIR_GLOBAL __attribute__((opencl_global)) @@ -111,4 +111,4 @@ static inline void atomicStore(int *Ptr, int V) { __spv::MemorySemanticsMask::SequentiallyConsistent, V); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/bfloat16_wrapper.cpp b/libdevice/bfloat16_wrapper.cpp index b2b8709f9dfbc..a0b6b96d4a293 100644 --- a/libdevice/bfloat16_wrapper.cpp +++ b/libdevice/bfloat16_wrapper.cpp @@ -8,7 +8,7 @@ #include "device.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #include #include @@ -23,4 +23,4 @@ float __devicelib_ConvertBF16ToFINTEL(const uint16_t &x) { return __spirv_ConvertBF16ToFINTEL(x); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/cmath_wrapper.cpp b/libdevice/cmath_wrapper.cpp index 5d9c8f0a77d13..c99ca6a889633 100644 --- a/libdevice/cmath_wrapper.cpp +++ b/libdevice/cmath_wrapper.cpp @@ -8,7 +8,7 @@ #include "device_math.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) DEVICE_EXTERN_C_INLINE int abs(int x) { return __devicelib_abs(x); } @@ -167,4 +167,4 @@ DEVICE_EXTERN_C_INLINE float rintf(float x) { return __nv_rintf(x); } #endif // __NVPTX__ -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ diff --git a/libdevice/cmath_wrapper_fp64.cpp b/libdevice/cmath_wrapper_fp64.cpp index e7b0815ae6526..d827b953520a2 100644 --- a/libdevice/cmath_wrapper_fp64.cpp +++ b/libdevice/cmath_wrapper_fp64.cpp @@ -9,7 +9,7 @@ #include "device_math.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) // All exported functions in math and complex device libraries are weak // reference. If users provide their own math or complex functions(with @@ -464,4 +464,4 @@ double _Sinh(double x, double y) { // compute y * sinh(x), |y| <= 1 } } #endif // defined(_WIN32) -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ diff --git a/libdevice/complex_wrapper.cpp b/libdevice/complex_wrapper.cpp index 552abfbf85bde..f01872fb6aee1 100644 --- a/libdevice/complex_wrapper.cpp +++ b/libdevice/complex_wrapper.cpp @@ -8,7 +8,7 @@ #include "device_complex.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) DEVICE_EXTERN_C_INLINE float cimagf(float __complex__ z) { return __devicelib_cimagf(z); } @@ -99,4 +99,4 @@ DEVICE_EXTERN_C_INLINE float __complex__ __divsc3(float __a, float __b, float __c, float __d) { return __devicelib___divsc3(__a, __b, __c, __d); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/complex_wrapper_fp64.cpp b/libdevice/complex_wrapper_fp64.cpp index 300eb107ca6d0..d3f42face5ec3 100644 --- a/libdevice/complex_wrapper_fp64.cpp +++ b/libdevice/complex_wrapper_fp64.cpp @@ -9,7 +9,7 @@ #include "device_complex.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) DEVICE_EXTERN_C_INLINE double cimag(double __complex__ z) { return __devicelib_cimag(z); } @@ -100,4 +100,4 @@ DEVICE_EXTERN_C_INLINE double __complex__ __divdc3(double __a, double __b, double __c, double __d) { return __devicelib___divdc3(__a, __b, __c, __d); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/crt_wrapper.cpp b/libdevice/crt_wrapper.cpp index 94481bc640de1..c9f21ed03975a 100644 --- a/libdevice/crt_wrapper.cpp +++ b/libdevice/crt_wrapper.cpp @@ -8,7 +8,7 @@ #include "wrapper.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) DEVICE_EXTERN_C_INLINE void *memcpy(void *dest, const void *src, size_t n) { return __devicelib_memcpy(dest, src, n); @@ -64,4 +64,4 @@ void __assert_fail(const char *expr, const char *file, unsigned int line, __spirv_LocalInvocationId_z()); } #endif -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ diff --git a/libdevice/device.h b/libdevice/device.h index 9702b52f1b391..360af54f9b4c4 100644 --- a/libdevice/device.h +++ b/libdevice/device.h @@ -15,7 +15,7 @@ #define EXTERN_C #endif // __cplusplus -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) #ifdef __SYCL_DEVICE_ONLY__ #define DEVICE_EXTERNAL SYCL_EXTERNAL __attribute__((weak)) #else // __SYCL_DEVICE_ONLY__ @@ -27,11 +27,11 @@ DEVICE_EXTERNAL EXTERN_C __attribute__((always_inline)) #define DEVICE_EXTERN_C_NOINLINE \ DEVICE_EXTERNAL EXTERN_C __attribute__((noinline)) -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ -#if defined(__SPIR__) || defined(__LIBDEVICE_HOST_IMPL__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__LIBDEVICE_HOST_IMPL__) #define __LIBDEVICE_IMF_ENABLED__ -#endif // __SPIR__ || __LIBDEVICE_HOST_IMPL__ +#endif // __SPIR__ || __SPIRV__ || __LIBDEVICE_HOST_IMPL__ #ifdef __LIBDEVICE_HOST_IMPL__ // For host implementation, all functions will be located in a static library diff --git a/libdevice/device_complex.h b/libdevice/device_complex.h index 317b631f85ac1..6fa4254ff2025 100644 --- a/libdevice/device_complex.h +++ b/libdevice/device_complex.h @@ -10,7 +10,7 @@ #include "device.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) // TODO: This needs to be more robust. // clang doesn't recognize the c11 CMPLX macro, but it does have @@ -165,5 +165,5 @@ double __complex__ __devicelib___divdc3(double a, double b, double c, double d); DEVICE_EXTERN_C float __complex__ __devicelib___divsc3(float a, float b, float c, float d); -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ #endif // __LIBDEVICE_DEVICE_COMPLEX_H_ diff --git a/libdevice/device_imf.hpp b/libdevice/device_imf.hpp index 9c55dfb7ad8d4..9580da09a9d0e 100644 --- a/libdevice/device_imf.hpp +++ b/libdevice/device_imf.hpp @@ -18,9 +18,10 @@ #include #ifdef __LIBDEVICE_IMF_ENABLED__ -#if !defined(__SPIR__) && !defined(__LIBDEVICE_HOST_IMPL__) +#if !defined(__SPIR__) && !defined(__SPIRV__) && \ + !defined(__LIBDEVICE_HOST_IMPL__) #error \ - "__SPIR__ or __LIBDEVICE_HOST_IMPL__ must be defined to enable device imf functions!" + "__SPIR__ or __SPIRV__ or __LIBDEVICE_HOST_IMPL__ must be defined to enable device imf functions!" #endif // TODO: Bitcast is valid to trivially copyable object only but using @@ -110,7 +111,7 @@ template static inline Ty __imin(Ty x, Ty y) { static inline float __fclamp(float x, float y, float z) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fmin(__builtin_fmax(x, y), z); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fclamp(x, y, z); #endif } @@ -119,7 +120,7 @@ static inline float __fclamp(float x, float y, float z) { static inline float __fma(float x, float y, float z) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fmaf(x, y, z); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fma(x, y, z); #endif } @@ -127,7 +128,7 @@ static inline float __fma(float x, float y, float z) { static inline double __fma(double x, double y, double z) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fma(x, y, z); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fma(x, y, z); #endif } @@ -142,7 +143,7 @@ static inline _iml_half __fma(_iml_half x, _iml_half y, _iml_half z) { float tmp_z = __half2float(z_i); float res = __builtin_fmaf(tmp_x, tmp_y, tmp_z); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_fma(x_i, y_i, z_i)); #endif } @@ -160,7 +161,7 @@ static inline _iml_bf16 __fma(_iml_bf16 x, _iml_bf16 y, _iml_bf16 z) { static inline float __sqrt(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_sqrtf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_sqrt(x); #endif } @@ -168,7 +169,7 @@ static inline float __sqrt(float x) { static inline double __sqrt(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_sqrt(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_sqrt(x); #endif } @@ -179,7 +180,7 @@ static inline _iml_half __sqrt(_iml_half x) { float tmp_x = __half2float(x_i); float res = __builtin_sqrtf(tmp_x); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_sqrt(x_i)); #endif } @@ -194,7 +195,7 @@ static inline _iml_bf16 __sqrt(_iml_bf16 x) { static inline float __rsqrt(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return 1.f / __builtin_sqrtf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_rsqrt(x); #endif } @@ -202,7 +203,7 @@ static inline float __rsqrt(float x) { static inline double __rsqrt(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return 1.0 / __builtin_sqrt(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_rsqrt(x); #endif } @@ -213,7 +214,7 @@ static inline _iml_half __rsqrt(_iml_half x) { float tmp_x = __half2float(x_i); float res = 1.f / __builtin_sqrtf(tmp_x); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_rsqrt(x_i)); #endif } @@ -228,7 +229,7 @@ static inline _iml_bf16 __rsqrt(_iml_bf16 x) { static inline float __fmin(float x, float y) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fminf(x, y); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fmin(x, y); #endif } @@ -236,7 +237,7 @@ static inline float __fmin(float x, float y) { static inline double __fmin(double x, double y) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fmin(x, y); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fmin(x, y); #endif } @@ -249,7 +250,7 @@ static inline _iml_half __fmin(_iml_half x, _iml_half y) { float tmp_y = __half2float(y_i); float res = __builtin_fminf(tmp_x, tmp_y); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_fmin(x_i, y_i)); #endif } @@ -265,7 +266,7 @@ static inline _iml_bf16 __fmin(_iml_bf16 x, _iml_bf16 y) { static inline float __fmax(float x, float y) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fmaxf(x, y); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fmax(x, y); #endif } @@ -273,7 +274,7 @@ static inline float __fmax(float x, float y) { static inline double __fmax(double x, double y) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fmax(x, y); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fmax(x, y); #endif } @@ -286,7 +287,7 @@ static inline _iml_half __fmax(_iml_half x, _iml_half y) { float tmp_y = __half2float(y_i); float res = __builtin_fmaxf(tmp_x, tmp_y); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_fmax(x_i, y_i)); #endif } @@ -302,7 +303,7 @@ static inline _iml_bf16 __fmax(_iml_bf16 x, _iml_bf16 y) { static inline float __copysign(float x, float y) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_copysignf(x, y); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_copysign(x, y); #endif } @@ -310,7 +311,7 @@ static inline float __copysign(float x, float y) { static inline double __copysign(double x, double y) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_copysign(x, y); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_copysign(x, y); #endif } @@ -323,7 +324,7 @@ static inline _iml_half __copysign(_iml_half x, _iml_half y) { float tmp_y = __half2float(y_i); float res = __builtin_copysignf(tmp_x, tmp_y); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_copysign(x_i, y_i)); #endif } @@ -339,7 +340,7 @@ static inline _iml_bf16 __copysign(_iml_bf16 x, _iml_bf16 y) { static inline float __fabs(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fabsf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fabs(x); #endif } @@ -347,7 +348,7 @@ static inline float __fabs(float x) { static inline double __fabs(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_fabs(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_fabs(x); #endif } @@ -358,7 +359,7 @@ static inline _iml_half __fabs(_iml_half x) { float tmp_x = __half2float(x_i); float res = __builtin_fabsf(tmp_x); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_fabs(x_i)); #endif } @@ -373,7 +374,7 @@ static inline _iml_bf16 __fabs(_iml_bf16 x) { static inline float __rint(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_rintf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_rint(x); #endif } @@ -381,7 +382,7 @@ static inline float __rint(float x) { static inline double __rint(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_rint(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_rint(x); #endif } @@ -392,7 +393,7 @@ static inline _iml_half __rint(_iml_half x) { float tmp_x = __half2float(x_i); float res = __builtin_rintf(tmp_x); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_rint(x_i)); #endif } @@ -407,7 +408,7 @@ static inline _iml_bf16 __rint(_iml_bf16 x) { static inline float __floor(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_floorf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_floor(x); #endif } @@ -415,7 +416,7 @@ static inline float __floor(float x) { static inline double __floor(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_floor(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_floor(x); #endif } @@ -426,7 +427,7 @@ static inline _iml_half __floor(_iml_half x) { float tmp_x = __half2float(x_i); float res = __builtin_floorf(tmp_x); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_floor(x_i)); #endif } @@ -441,7 +442,7 @@ static inline _iml_bf16 __floor(_iml_bf16 x) { static inline float __ceil(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_ceilf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_ceil(x); #endif } @@ -449,7 +450,7 @@ static inline float __ceil(float x) { static inline double __ceil(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_ceil(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_ceil(x); #endif } @@ -460,7 +461,7 @@ static inline _iml_half __ceil(_iml_half x) { float tmp_x = __half2float(x_i); float res = __builtin_ceilf(tmp_x); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_ceil(x_i)); #endif } @@ -475,7 +476,7 @@ static inline _iml_bf16 __ceil(_iml_bf16 x) { static inline float __trunc(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_truncf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_trunc(x); #endif } @@ -483,7 +484,7 @@ static inline float __trunc(float x) { static inline double __trunc(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_trunc(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_trunc(x); #endif } @@ -491,7 +492,7 @@ static inline double __trunc(double x) { static inline float __fast_exp10f(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_expf(0x1.26bb1cp1f * x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_native_exp(0x1.26bb1cp1f * x); #endif } @@ -499,7 +500,7 @@ static inline float __fast_exp10f(float x) { static inline float __fast_expf(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_expf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_native_exp(x); #endif } @@ -507,7 +508,7 @@ static inline float __fast_expf(float x) { static inline float __fast_logf(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_logf(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_native_log(x); #endif } @@ -515,7 +516,7 @@ static inline float __fast_logf(float x) { static inline float __fast_log2f(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_log2f(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_native_log(x) / 0x1.62e43p-1f; #endif } @@ -523,7 +524,7 @@ static inline float __fast_log2f(float x) { static inline float __fast_log10f(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_log10f(x); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_native_log(x) / 0x1.26bb1cp1f; #endif } @@ -531,7 +532,7 @@ static inline float __fast_log10f(float x) { static inline float __fast_powf(float x, float y) { #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_powf(x, y); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_native_powr(x, y); #endif } @@ -554,7 +555,7 @@ static inline float __fast_fdividef(float x, float y) { #if defined(__LIBDEVICE_HOST_IMPL__) return x / y; -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_native_divide(x, y); #endif } @@ -565,7 +566,7 @@ static inline _iml_half __trunc(_iml_half x) { float tmp_x = __half2float(x_i); float res = __builtin_truncf(tmp_x); return _iml_half(__float2half(res)); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return _iml_half(__spirv_ocl_trunc(x_i)); #endif } @@ -582,7 +583,7 @@ static inline int __clz(int x) { uint32_t xi32 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_clz(xi32); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_clz(xi32); #endif } @@ -593,7 +594,7 @@ static inline int __clzll(long long int x) { uint64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_clzll(xi64); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_clz(xi64); #endif } @@ -602,7 +603,7 @@ static inline int __popc(unsigned int x) { uint32_t xui32 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_popcount(xui32); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_popcount(xui32); #endif } @@ -611,7 +612,7 @@ static inline int __popcll(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __builtin_popcountll(xui64); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_popcount(xui64); #endif } @@ -665,7 +666,7 @@ template static inline Ty __uhadd(Ty x, Ty y) { "__uhadd can only accept unsigned integral type."); #if defined(__LIBDEVICE_HOST_IMPL__) return (x >> 1) + (y >> 1) + ((x & y) & 0x1); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_u_hadd(x, y); #endif } @@ -675,7 +676,7 @@ template static inline Ty __shadd(Ty x, Ty y) { "__shadd can only accept signed integral type."); #if defined(__LIBDEVICE_HOST_IMPL__) return (x >> 1) + (y >> 1) + ((x & y) & 0x1); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_s_hadd(x, y); #endif } @@ -685,7 +686,7 @@ template static inline Ty __urhadd(Ty x, Ty y) { "__urhadd can only accept unsigned integral type."); #if defined(__LIBDEVICE_HOST_IMPL__) return (x >> 1) + (y >> 1) + ((x | y) & 0x1); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_u_rhadd(x, y); #endif } @@ -695,7 +696,7 @@ template static inline Ty __srhadd(Ty x, Ty y) { "__srhadd can only accept signed integral type."); #if defined(__LIBDEVICE_HOST_IMPL__) return (x >> 1) + (y >> 1) + ((x | y) & 0x1); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_s_rhadd(x, y); #endif } diff --git a/libdevice/device_itt.h b/libdevice/device_itt.h index 77fdf01939432..e5513d04709ef 100644 --- a/libdevice/device_itt.h +++ b/libdevice/device_itt.h @@ -11,7 +11,7 @@ #include "device.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #include "spirv_vars.h" #define ITT_STUB_ATTRIBUTES __attribute__((noinline, optnone)) @@ -108,5 +108,5 @@ SYCL_EXTERNAL EXTERN_C void __itt_offload_atomic_op_finish(void *object, __itt_atomic_mem_op_t op_type, __itt_atomic_mem_order_t mem_order); -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ #endif // __LIBDEVICE_DEVICE_ITT_H__ diff --git a/libdevice/device_math.h b/libdevice/device_math.h index a402c748299d2..650f4aefbc942 100644 --- a/libdevice/device_math.h +++ b/libdevice/device_math.h @@ -10,7 +10,7 @@ #define __LIBDEVICE_DEVICE_MATH_H__ #include "device.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) #include typedef struct { @@ -307,5 +307,5 @@ float __devicelib_scalbnf(float x, int n); DEVICE_EXTERN_C double __devicelib_scalbn(double x, int exp); -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ #endif // __LIBDEVICE_DEVICE_MATH_H__ diff --git a/libdevice/fallback-bfloat16.cpp b/libdevice/fallback-bfloat16.cpp index e5596ff4871dd..84015d03b35b0 100644 --- a/libdevice/fallback-bfloat16.cpp +++ b/libdevice/fallback-bfloat16.cpp @@ -8,7 +8,7 @@ #include "device.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #include @@ -43,4 +43,4 @@ __devicelib_ConvertBF16ToFINTEL(const uint16_t &a) { return floatValue; } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/fallback-cassert.cpp b/libdevice/fallback-cassert.cpp index 47bae9f54714a..5d3c99d63c556 100644 --- a/libdevice/fallback-cassert.cpp +++ b/libdevice/fallback-cassert.cpp @@ -10,7 +10,7 @@ #include "include/assert-happened.hpp" #include "wrapper.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #define ASSERT_NONE 0 #define ASSERT_START 1 @@ -98,7 +98,7 @@ DEVICE_EXTERN_C void __devicelib_assert_fail(const char *expr, const char *file, // volatile int *die = (int *)0x0; // *die = 0xdead; } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ #ifdef __NVPTX__ diff --git a/libdevice/fallback-cmath-fp64.cpp b/libdevice/fallback-cmath-fp64.cpp index 9656f229c4fd1..9111674fa9d2c 100644 --- a/libdevice/fallback-cmath-fp64.cpp +++ b/libdevice/fallback-cmath-fp64.cpp @@ -9,7 +9,7 @@ #include "device_math.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) // To support fallback device libraries on-demand loading, please update the // DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add @@ -159,4 +159,4 @@ DEVICE_EXTERN_C_INLINE double __devicelib_scalbn(double x, int exp) { return __spirv_ocl_ldexp(x, exp); } -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ diff --git a/libdevice/fallback-cmath.cpp b/libdevice/fallback-cmath.cpp index dc9e2806111f5..cd660e59a54e9 100644 --- a/libdevice/fallback-cmath.cpp +++ b/libdevice/fallback-cmath.cpp @@ -8,7 +8,7 @@ #include "device_math.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) // To support fallback device libraries on-demand loading, please update the // DeviceLibFuncMap in llvm/tools/sycl-post-link/sycl-post-link.cpp if you add @@ -175,4 +175,4 @@ float __devicelib_asinhf(float x) { return __spirv_ocl_asinh(x); } DEVICE_EXTERN_C_INLINE float __devicelib_atanhf(float x) { return __spirv_ocl_atanh(x); } -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ diff --git a/libdevice/fallback-complex-fp64.cpp b/libdevice/fallback-complex-fp64.cpp index 95f7734524ac6..5ca69c0100962 100644 --- a/libdevice/fallback-complex-fp64.cpp +++ b/libdevice/fallback-complex-fp64.cpp @@ -9,7 +9,7 @@ #include "device_complex.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #include // To support fallback device libraries on-demand loading, please update the @@ -428,4 +428,4 @@ double __complex__ __devicelib_catan(double __complex__ z) { __devicelib_catanh(CMPLX(-__devicelib_cimag(z), __devicelib_creal(z))); return CMPLX(__devicelib_cimag(w), -__devicelib_creal(w)); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/fallback-complex.cpp b/libdevice/fallback-complex.cpp index 27b1d5c0dae94..daa8c234fbc88 100644 --- a/libdevice/fallback-complex.cpp +++ b/libdevice/fallback-complex.cpp @@ -8,7 +8,7 @@ #include "device_complex.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #include // To support fallback device libraries on-demand loading, please update the @@ -431,4 +431,4 @@ float __complex__ __devicelib_catanf(float __complex__ z) { CMPLXF(-__devicelib_cimagf(z), __devicelib_crealf(z))); return CMPLXF(__devicelib_cimagf(w), -__devicelib_crealf(w)); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/fallback-cstring.cpp b/libdevice/fallback-cstring.cpp index bebfc621857d7..5d384f00a78cb 100644 --- a/libdevice/fallback-cstring.cpp +++ b/libdevice/fallback-cstring.cpp @@ -9,7 +9,7 @@ #include "wrapper.h" #include -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) static void *__devicelib_memcpy_uint8_aligned(void *dest, const void *src, size_t n) { @@ -202,4 +202,4 @@ int __devicelib_memcmp(const void *s1, const void *s2, size_t n) { return head_cmp; } -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ diff --git a/libdevice/imf_half.hpp b/libdevice/imf_half.hpp index 027cdd7a3e386..076cfdf731561 100644 --- a/libdevice/imf_half.hpp +++ b/libdevice/imf_half.hpp @@ -17,7 +17,7 @@ #include #ifdef __LIBDEVICE_IMF_ENABLED__ -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) typedef _Float16 _iml_half_internal; #else typedef uint16_t _iml_half_internal; @@ -426,7 +426,7 @@ static uint16_t __iml_integral2half_s(Ty i, __iml_rounding_mode rounding_mode) { static inline _iml_half_internal __float2half(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __iml_fp2half(x, __IML_RTE); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rhalf_rte(x); #endif } @@ -469,7 +469,7 @@ static inline float __half2float(_iml_half_internal x) { fp32_bits |= (exp32 << 23); fp32_bits |= frac32; return __builtin_bit_cast(float, fp32_bits); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rfloat_rte(x); #endif } @@ -492,7 +492,7 @@ class _iml_half { return _half_internal == rh._half_internal; } bool operator!=(const _iml_half &rh) { return !operator==(rh); } -#if (__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) _iml_half &operator+=(const _iml_half &rh) { _half_internal += rh._half_internal; return *this; diff --git a/libdevice/imf_utils/double_convert.cpp b/libdevice/imf_utils/double_convert.cpp index 0cfc7ef693a02..c4cd6dea07bf1 100644 --- a/libdevice/imf_utils/double_convert.cpp +++ b/libdevice/imf_utils/double_convert.cpp @@ -14,7 +14,7 @@ static inline float __double2float_rd(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rfloat_rtn(x); #endif } @@ -22,7 +22,7 @@ static inline float __double2float_rd(double x) { static inline float __double2float_rn(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rfloat_rte(x); #endif } @@ -30,7 +30,7 @@ static inline float __double2float_rn(double x) { static inline float __double2float_ru(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rfloat_rtp(x); #endif } @@ -38,7 +38,7 @@ static inline float __double2float_ru(double x) { static inline float __double2float_rz(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rfloat_rtz(x); #endif } @@ -58,7 +58,7 @@ float __devicelib_imf_double2float_rz(double x) { return __double2float_rz(x); } static inline int __double2int_rd(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rtn(x); #endif } @@ -66,7 +66,7 @@ static inline int __double2int_rd(double x) { static inline int __double2int_rn(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rte(x); #endif } @@ -74,7 +74,7 @@ static inline int __double2int_rn(double x) { static inline int __double2int_ru(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rtp(x); #endif } @@ -82,7 +82,7 @@ static inline int __double2int_ru(double x) { static inline int __double2int_rz(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rtz(x); #endif } @@ -119,7 +119,7 @@ static inline unsigned int __double2uint_rd(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rtn(x); #endif } @@ -129,7 +129,7 @@ static inline unsigned int __double2uint_rn(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rte(x); #endif } @@ -139,7 +139,7 @@ static inline unsigned int __double2uint_ru(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rtp(x); #endif } @@ -149,7 +149,7 @@ static inline unsigned int __double2uint_rz(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rtz(x); #endif } @@ -177,7 +177,7 @@ unsigned int __devicelib_imf_double2uint_rz(double x) { static inline long long int __double2ll_rd(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rtn(x); #endif } @@ -185,7 +185,7 @@ static inline long long int __double2ll_rd(double x) { static inline long long int __double2ll_rn(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rte(x); #endif } @@ -193,7 +193,7 @@ static inline long long int __double2ll_rn(double x) { static inline long long int __double2ll_ru(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rtp(x); #endif } @@ -201,7 +201,7 @@ static inline long long int __double2ll_ru(double x) { static inline long long int __double2ll_rz(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rtz(x); #endif } @@ -231,7 +231,7 @@ static inline unsigned long long int __double2ull_rd(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rtn(x); #endif } @@ -241,7 +241,7 @@ static inline unsigned long long int __double2ull_rn(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rte(x); #endif } @@ -251,7 +251,7 @@ static inline unsigned long long int __double2ull_ru(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rtp(x); #endif } @@ -261,7 +261,7 @@ static inline unsigned long long int __double2ull_rz(double x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __double2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rtz(x); #endif } @@ -304,7 +304,7 @@ double __devicelib_imf_hiloint2double(int hi, int lo) { static inline double __int2double_rn(int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rdouble(x); #endif } @@ -316,7 +316,7 @@ static inline double __ll2double_rd(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rdouble_rtn(xi64); #endif } @@ -325,7 +325,7 @@ static inline double __ll2double_rn(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rdouble_rte(xi64); #endif } @@ -334,7 +334,7 @@ static inline double __ll2double_ru(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rdouble_rtp(xi64); #endif } @@ -343,7 +343,7 @@ static inline double __ll2double_rz(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rdouble_rtz(xi64); #endif } @@ -376,7 +376,7 @@ double __devicelib_imf_longlong_as_double(long long int x) { static inline double __uint2double_rn(unsigned int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rdouble_rte(x); #endif } @@ -390,7 +390,7 @@ static inline double __ull2double_rd(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rdouble_rtn(xui64); #endif } @@ -399,7 +399,7 @@ static inline double __ull2double_rn(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rdouble_rte(xui64); #endif } @@ -408,7 +408,7 @@ static inline double __ull2double_ru(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rdouble_rtp(xui64); #endif } @@ -417,7 +417,7 @@ static inline double __ull2double_rz(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rdouble_rtz(xui64); #endif } @@ -446,7 +446,7 @@ DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_double2half(double x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __iml_fp2half(x, __IML_RTE); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rhalf_rte(x); #endif } diff --git a/libdevice/imf_utils/float_convert.cpp b/libdevice/imf_utils/float_convert.cpp index d42749c15b0aa..85299c0f33823 100644 --- a/libdevice/imf_utils/float_convert.cpp +++ b/libdevice/imf_utils/float_convert.cpp @@ -14,7 +14,7 @@ static inline int __float2int_rd(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rtn(x); #endif } @@ -22,7 +22,7 @@ static inline int __float2int_rd(float x) { static inline int __float2int_rn(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rte(x); #endif } @@ -30,7 +30,7 @@ static inline int __float2int_rn(float x) { static inline int __float2int_ru(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rtp(x); #endif } @@ -38,7 +38,7 @@ static inline int __float2int_ru(float x) { static inline int __float2int_rz(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_rtz(x); #endif } @@ -60,7 +60,7 @@ static inline unsigned int __float2uint_rd(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rtn(x); #endif } @@ -70,7 +70,7 @@ static inline unsigned int __float2uint_rn(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rte(x); #endif } @@ -80,7 +80,7 @@ static inline unsigned int __float2uint_ru(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rtp(x); #endif } @@ -90,7 +90,7 @@ static inline unsigned int __float2uint_rz(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_rtz(x); #endif } @@ -118,7 +118,7 @@ unsigned int __devicelib_imf_float2uint_rz(float x) { static inline long long int __float2ll_rd(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rtn(x); #endif } @@ -126,7 +126,7 @@ static inline long long int __float2ll_rd(float x) { static inline long long int __float2ll_rn(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rte(x); #endif } @@ -134,7 +134,7 @@ static inline long long int __float2ll_rn(float x) { static inline long long int __float2ll_ru(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rtp(x); #endif } @@ -142,7 +142,7 @@ static inline long long int __float2ll_ru(float x) { static inline long long int __float2ll_rz(float x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_rtz(x); #endif } @@ -164,7 +164,7 @@ static inline unsigned long long int __float2ull_rd(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rtn(x); #endif } @@ -174,7 +174,7 @@ static inline unsigned long long int __float2ull_rn(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rte(x); #endif } @@ -184,7 +184,7 @@ static inline unsigned long long int __float2ull_ru(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rtp(x); #endif } @@ -194,7 +194,7 @@ static inline unsigned long long int __float2ull_rz(float x) { return 0; #if defined(__LIBDEVICE_HOST_IMPL__) return __float2Tp_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_rtz(x); #endif } @@ -230,7 +230,7 @@ unsigned int __devicelib_imf_float_as_uint(float x) { static inline float __int2float_rd(int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rtn(x); #endif } @@ -238,7 +238,7 @@ static inline float __int2float_rd(int x) { static inline float __int2float_rn(int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rte(x); #endif } @@ -246,7 +246,7 @@ static inline float __int2float_rn(int x) { static inline float __int2float_ru(int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rtp(x); #endif } @@ -254,7 +254,7 @@ static inline float __int2float_ru(int x) { static inline float __int2float_rz(int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rtz(x); #endif } @@ -278,7 +278,7 @@ static inline float __ll2float_rd(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rtn(xi64); #endif } @@ -287,7 +287,7 @@ static inline float __ll2float_rn(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rte(xi64); #endif } @@ -296,7 +296,7 @@ static inline float __ll2float_ru(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rtp(xi64); #endif } @@ -305,7 +305,7 @@ static inline float __ll2float_rz(long long int x) { int64_t xi64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xi64, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rfloat_rtz(xi64); #endif } @@ -325,7 +325,7 @@ float __devicelib_imf_ll2float_rz(long long int x) { return __ll2float_rz(x); } static inline float __uint2float_rd(unsigned int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rtn(x); #endif } @@ -333,7 +333,7 @@ static inline float __uint2float_rd(unsigned int x) { static inline float __uint2float_rn(unsigned int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rte(x); #endif } @@ -341,7 +341,7 @@ static inline float __uint2float_rn(unsigned int x) { static inline float __uint2float_ru(unsigned int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rtp(x); #endif } @@ -349,7 +349,7 @@ static inline float __uint2float_ru(unsigned int x) { static inline float __uint2float_rz(unsigned int x) { #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(x, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rtz(x); #endif } @@ -384,7 +384,7 @@ float __devicelib_imf_ull2float_rd(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_DOWNWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rtn(xui64); #endif } @@ -394,7 +394,7 @@ float __devicelib_imf_ull2float_rn(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_TONEAREST); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rte(xui64); #endif } @@ -404,7 +404,7 @@ float __devicelib_imf_ull2float_ru(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_UPWARD); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rtp(xui64); #endif } @@ -414,7 +414,7 @@ float __devicelib_imf_ull2float_rz(unsigned long long int x) { uint64_t xui64 = x; #if defined(__LIBDEVICE_HOST_IMPL__) return __integral2FP_host(xui64, FE_TOWARDZERO); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rfloat_rtz(xui64); #endif } diff --git a/libdevice/imf_utils/half_convert.cpp b/libdevice/imf_utils/half_convert.cpp index 80bf2ea064da0..3e23d3a46f01e 100644 --- a/libdevice/imf_utils/half_convert.cpp +++ b/libdevice/imf_utils/half_convert.cpp @@ -18,7 +18,7 @@ float __devicelib_imf_half2float(_iml_half_internal x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_float2half_rd(float x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rhalf_rtn(x); #else return __iml_fp2half(x, __IML_RTN); @@ -27,7 +27,7 @@ _iml_half_internal __devicelib_imf_float2half_rd(float x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_float2half_rn(float x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rhalf_rte(x); #else return __iml_fp2half(x, __IML_RTE); @@ -36,7 +36,7 @@ _iml_half_internal __devicelib_imf_float2half_rn(float x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_float2half_ru(float x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rhalf_rtp(x); #else return __iml_fp2half(x, __IML_RTP); @@ -45,7 +45,7 @@ _iml_half_internal __devicelib_imf_float2half_ru(float x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_float2half_rz(float x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_FConvert_Rhalf_rtz(x); #else return __iml_fp2half(x, __IML_RTZ); @@ -54,7 +54,7 @@ _iml_half_internal __devicelib_imf_float2half_rz(float x) { DEVICE_EXTERN_C_INLINE int __devicelib_imf_half2int_rd(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_sat_rtn(h); #else return __iml_half2integral_s(h, __IML_RTN); @@ -63,7 +63,7 @@ int __devicelib_imf_half2int_rd(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE int __devicelib_imf_half2int_rn(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_sat_rte(h); #else return __iml_half2integral_s(h, __IML_RTE); @@ -72,7 +72,7 @@ int __devicelib_imf_half2int_rn(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE int __devicelib_imf_half2int_ru(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_sat_rtp(h); #else return __iml_half2integral_s(h, __IML_RTP); @@ -81,7 +81,7 @@ int __devicelib_imf_half2int_ru(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE int __devicelib_imf_half2int_rz(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rint_sat_rtz(h); #else return __iml_half2integral_s(h, __IML_RTZ); @@ -90,7 +90,7 @@ int __devicelib_imf_half2int_rz(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE long long __devicelib_imf_half2ll_rd(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_sat_rtn(h); #else return __iml_half2integral_s(h, __IML_RTN); @@ -99,7 +99,7 @@ long long __devicelib_imf_half2ll_rd(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE long long __devicelib_imf_half2ll_rn(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_sat_rte(h); #else return __iml_half2integral_s(h, __IML_RTE); @@ -108,7 +108,7 @@ long long __devicelib_imf_half2ll_rn(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE long long __devicelib_imf_half2ll_ru(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_sat_rtp(h); #else return __iml_half2integral_s(h, __IML_RTP); @@ -117,7 +117,7 @@ long long __devicelib_imf_half2ll_ru(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE long long __devicelib_imf_half2ll_rz(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rlong_sat_rtz(h); #else return __iml_half2integral_s(h, __IML_RTZ); @@ -126,7 +126,7 @@ long long __devicelib_imf_half2ll_rz(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE short __devicelib_imf_half2short_rd(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rshort_sat_rtn(h); #else return __iml_half2integral_s(h, __IML_RTN); @@ -135,7 +135,7 @@ short __devicelib_imf_half2short_rd(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE short __devicelib_imf_half2short_rn(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rshort_sat_rte(h); #else return __iml_half2integral_s(h, __IML_RTE); @@ -144,7 +144,7 @@ short __devicelib_imf_half2short_rn(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE short __devicelib_imf_half2short_ru(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rshort_sat_rtp(h); #else return __iml_half2integral_s(h, __IML_RTP); @@ -153,7 +153,7 @@ short __devicelib_imf_half2short_ru(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE short __devicelib_imf_half2short_rz(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToS_Rshort_sat_rtz(h); #else return __iml_half2integral_s(h, __IML_RTZ); @@ -162,7 +162,7 @@ short __devicelib_imf_half2short_rz(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned int __devicelib_imf_half2uint_rd(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_sat_rtn(h); #else return __iml_half2integral_u(h, __IML_RTN); @@ -171,7 +171,7 @@ unsigned int __devicelib_imf_half2uint_rd(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned int __devicelib_imf_half2uint_rn(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_sat_rte(h); #else return __iml_half2integral_u(h, __IML_RTE); @@ -180,7 +180,7 @@ unsigned int __devicelib_imf_half2uint_rn(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned int __devicelib_imf_half2uint_ru(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_sat_rtp(h); #else return __iml_half2integral_u(h, __IML_RTP); @@ -189,7 +189,7 @@ unsigned int __devicelib_imf_half2uint_ru(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned int __devicelib_imf_half2uint_rz(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Ruint_sat_rtz(h); #else return __iml_half2integral_u(h, __IML_RTZ); @@ -198,7 +198,7 @@ unsigned int __devicelib_imf_half2uint_rz(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned long long __devicelib_imf_half2ull_rd(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_sat_rtn(h); #else return __iml_half2integral_u(h, __IML_RTN); @@ -207,7 +207,7 @@ unsigned long long __devicelib_imf_half2ull_rd(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned long long __devicelib_imf_half2ull_rn(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_sat_rte(h); #else return __iml_half2integral_u(h, __IML_RTE); @@ -216,7 +216,7 @@ unsigned long long __devicelib_imf_half2ull_rn(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned long long __devicelib_imf_half2ull_ru(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_sat_rtp(h); #else return __iml_half2integral_u(h, __IML_RTP); @@ -225,7 +225,7 @@ unsigned long long __devicelib_imf_half2ull_ru(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned long long __devicelib_imf_half2ull_rz(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rulong_sat_rtz(h); #else return __iml_half2integral_u(h, __IML_RTZ); @@ -234,7 +234,7 @@ unsigned long long __devicelib_imf_half2ull_rz(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned short __devicelib_imf_half2ushort_rd(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rushort_sat_rtn(h); #else return __iml_half2integral_u(h, __IML_RTN); @@ -243,7 +243,7 @@ unsigned short __devicelib_imf_half2ushort_rd(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned short __devicelib_imf_half2ushort_rn(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rushort_sat_rte(h); #else return __iml_half2integral_u(h, __IML_RTE); @@ -252,7 +252,7 @@ unsigned short __devicelib_imf_half2ushort_rn(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned short __devicelib_imf_half2ushort_ru(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rushort_sat_rtp(h); #else return __iml_half2integral_u(h, __IML_RTP); @@ -261,7 +261,7 @@ unsigned short __devicelib_imf_half2ushort_ru(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE unsigned short __devicelib_imf_half2ushort_rz(_iml_half_internal h) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertFToU_Rushort_sat_rtz(h); #else return __iml_half2integral_u(h, __IML_RTZ); @@ -280,7 +280,7 @@ unsigned short __devicelib_imf_half_as_ushort(_iml_half_internal h) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_int2half_rd(int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtn(x); #else return __iml_integral2half_s(x, __IML_RTN); @@ -289,7 +289,7 @@ _iml_half_internal __devicelib_imf_int2half_rd(int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_int2half_rn(int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rte(x); #else return __iml_integral2half_s(x, __IML_RTE); @@ -298,7 +298,7 @@ _iml_half_internal __devicelib_imf_int2half_rn(int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_int2half_ru(int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtp(x); #else return __iml_integral2half_s(x, __IML_RTP); @@ -307,7 +307,7 @@ _iml_half_internal __devicelib_imf_int2half_ru(int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_int2half_rz(int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtz(x); #else return __iml_integral2half_s(x, __IML_RTZ); @@ -316,7 +316,7 @@ _iml_half_internal __devicelib_imf_int2half_rz(int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ll2half_rd(long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtn((int64_t)x); #else return __iml_integral2half_s(x, __IML_RTN); @@ -325,7 +325,7 @@ _iml_half_internal __devicelib_imf_ll2half_rd(long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ll2half_rn(long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rte((int64_t)x); #else return __iml_integral2half_s(x, __IML_RTE); @@ -334,7 +334,7 @@ _iml_half_internal __devicelib_imf_ll2half_rn(long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ll2half_ru(long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtp((int64_t)x); #else return __iml_integral2half_s(x, __IML_RTP); @@ -343,7 +343,7 @@ _iml_half_internal __devicelib_imf_ll2half_ru(long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ll2half_rz(long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtz((int64_t)x); #else return __iml_integral2half_s(x, __IML_RTZ); @@ -352,7 +352,7 @@ _iml_half_internal __devicelib_imf_ll2half_rz(long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_short2half_rd(short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtn(x); #else return __iml_integral2half_s(x, __IML_RTN); @@ -361,7 +361,7 @@ _iml_half_internal __devicelib_imf_short2half_rd(short x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_short2half_rn(short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rte(x); #else return __iml_integral2half_s(x, __IML_RTE); @@ -370,7 +370,7 @@ _iml_half_internal __devicelib_imf_short2half_rn(short x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_short2half_ru(short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtp(x); #else return __iml_integral2half_s(x, __IML_RTP); @@ -379,7 +379,7 @@ _iml_half_internal __devicelib_imf_short2half_ru(short x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_short2half_rz(short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertSToF_Rhalf_rtz(x); #else return __iml_integral2half_s(x, __IML_RTZ); @@ -393,7 +393,7 @@ _iml_half_internal __devicelib_imf_short_as_half(short x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_uint2half_rd(unsigned int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtn(x); #else return __iml_integral2half_u(x, __IML_RTN); @@ -402,7 +402,7 @@ _iml_half_internal __devicelib_imf_uint2half_rd(unsigned int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_uint2half_rn(unsigned int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rte(x); #else return __iml_integral2half_u(x, __IML_RTE); @@ -411,7 +411,7 @@ _iml_half_internal __devicelib_imf_uint2half_rn(unsigned int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_uint2half_ru(unsigned int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtp(x); #else return __iml_integral2half_u(x, __IML_RTP); @@ -420,7 +420,7 @@ _iml_half_internal __devicelib_imf_uint2half_ru(unsigned int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_uint2half_rz(unsigned int x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtz(x); #else return __iml_integral2half_u(x, __IML_RTZ); @@ -429,7 +429,7 @@ _iml_half_internal __devicelib_imf_uint2half_rz(unsigned int x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ull2half_rd(unsigned long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtn((uint64_t)x); #else return __iml_integral2half_u(x, __IML_RTN); @@ -438,7 +438,7 @@ _iml_half_internal __devicelib_imf_ull2half_rd(unsigned long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ull2half_rn(unsigned long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rte((uint64_t)x); #else return __iml_integral2half_u(x, __IML_RTE); @@ -447,7 +447,7 @@ _iml_half_internal __devicelib_imf_ull2half_rn(unsigned long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ull2half_ru(unsigned long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtp((uint64_t)x); #else return __iml_integral2half_u(x, __IML_RTP); @@ -456,7 +456,7 @@ _iml_half_internal __devicelib_imf_ull2half_ru(unsigned long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ull2half_rz(unsigned long long x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtz((uint64_t)x); #else return __iml_integral2half_u(x, __IML_RTZ); @@ -465,7 +465,7 @@ _iml_half_internal __devicelib_imf_ull2half_rz(unsigned long long x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ushort2half_rd(unsigned short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtn(x); #else return __iml_integral2half_u(x, __IML_RTN); @@ -474,7 +474,7 @@ _iml_half_internal __devicelib_imf_ushort2half_rd(unsigned short x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ushort2half_rn(unsigned short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rte(x); #else return __iml_integral2half_u(x, __IML_RTE); @@ -483,7 +483,7 @@ _iml_half_internal __devicelib_imf_ushort2half_rn(unsigned short x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ushort2half_ru(unsigned short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtp(x); #else return __iml_integral2half_u(x, __IML_RTP); @@ -492,7 +492,7 @@ _iml_half_internal __devicelib_imf_ushort2half_ru(unsigned short x) { DEVICE_EXTERN_C_INLINE _iml_half_internal __devicelib_imf_ushort2half_rz(unsigned short x) { -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) return __spirv_ConvertUToF_Rhalf_rtz(x); #else return __iml_integral2half_u(x, __IML_RTZ); diff --git a/libdevice/imf_utils/integer_misc.cpp b/libdevice/imf_utils/integer_misc.cpp index 1724f015bf658..fdc850ee42281 100644 --- a/libdevice/imf_utils/integer_misc.cpp +++ b/libdevice/imf_utils/integer_misc.cpp @@ -114,7 +114,7 @@ DEVICE_EXTERN_C_INLINE int __devicelib_imf_mul24(int x, int y) { #if defined(__LIBDEVICE_HOST_IMPL__) return x * y; -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_s_mul24(x, y); #endif } @@ -123,7 +123,7 @@ DEVICE_EXTERN_C_INLINE unsigned int __devicelib_imf_umul24(unsigned int x, unsigned int y) { #if defined(__LIBDEVICE_HOST_IMPL__) return x * y; -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_u_mul24(x, y); #endif } @@ -134,7 +134,7 @@ int __devicelib_imf_mulhi(int x, int y) { int64_t p = static_cast(x) * static_cast(y); p >>= 32; return static_cast(p); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_s_mul_hi(x, y); #endif } @@ -145,7 +145,7 @@ unsigned int __devicelib_imf_umulhi(unsigned int x, unsigned int y) { uint64_t p = static_cast(x) * static_cast(y); p >>= 32; return static_cast(p); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_u_mul_hi(x, y); #endif } @@ -156,7 +156,7 @@ long long int __devicelib_imf_mul64hi(long long int x, long long int y) { __int128_t p = static_cast<__int128_t>(x) * static_cast<__int128_t>(y); p >>= 64; return static_cast(p); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_s_mul_hi(static_cast(x), static_cast(y)); #endif } @@ -168,7 +168,7 @@ unsigned long long int __devicelib_imf_umul64hi(unsigned long long int x, __uint128_t p = static_cast<__uint128_t>(x) * static_cast<__uint128_t>(y); p >>= 64; return static_cast(p); -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) return __spirv_ocl_u_mul_hi(static_cast(x), static_cast(y)); #endif diff --git a/libdevice/include/assert-happened.hpp b/libdevice/include/assert-happened.hpp index 26e820c517d46..e6db239d8c315 100644 --- a/libdevice/include/assert-happened.hpp +++ b/libdevice/include/assert-happened.hpp @@ -10,7 +10,7 @@ // Treat this header as system one to workaround frontend's restriction #pragma clang system_header -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) // NOTE Layout of this structure should be aligned with the one in // sycl/include/sycl/detail/assert_happened.hpp diff --git a/libdevice/itt_compiler_wrappers.cpp b/libdevice/itt_compiler_wrappers.cpp index d1afd11fea4ef..c9ac0700fcbd7 100644 --- a/libdevice/itt_compiler_wrappers.cpp +++ b/libdevice/itt_compiler_wrappers.cpp @@ -8,7 +8,7 @@ #include "device_itt.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) SYCL_EXTERNAL EXTERN_C void __itt_offload_wi_start_wrapper() { if (!isITTEnabled()) @@ -53,4 +53,4 @@ SYCL_EXTERNAL EXTERN_C void __itt_offload_wi_resume_wrapper() { __itt_offload_wi_resume_stub(GroupID, WIID); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/itt_stubs.cpp b/libdevice/itt_stubs.cpp index 487f18cd0b5b2..62b0f27538ec2 100644 --- a/libdevice/itt_stubs.cpp +++ b/libdevice/itt_stubs.cpp @@ -8,7 +8,7 @@ #include "device_itt.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) SYCL_EXTERNAL EXTERN_C ITT_STUB_ATTRIBUTES void __itt_offload_wi_start_stub(size_t *group_id, size_t wi_id, uint32_t wg_size) {} @@ -35,4 +35,4 @@ SYCL_EXTERNAL EXTERN_C ITT_STUB_ATTRIBUTES void __itt_offload_atomic_op_finish_stub(void *object, __itt_atomic_mem_op_t op_type, __itt_atomic_mem_order_t mem_order) {} -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/itt_user_wrappers.cpp b/libdevice/itt_user_wrappers.cpp index 60206658df005..f2fe0ab7bfc74 100644 --- a/libdevice/itt_user_wrappers.cpp +++ b/libdevice/itt_user_wrappers.cpp @@ -8,7 +8,7 @@ #include "device_itt.h" -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) SYCL_EXTERNAL EXTERN_C void __itt_offload_wi_start(size_t *group_id, size_t wi_id, uint32_t wg_size) { @@ -63,4 +63,4 @@ __itt_offload_atomic_op_finish(void *object, __itt_atomic_mem_op_t op_type, __itt_offload_atomic_op_finish_stub(object, op_type, mem_order); } -#endif // __SPIR__ +#endif // __SPIR__ || __SPIRV__ diff --git a/libdevice/msvc_math.cpp b/libdevice/msvc_math.cpp index 38c4a99991a3e..3c7d8217fa9fe 100644 --- a/libdevice/msvc_math.cpp +++ b/libdevice/msvc_math.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -#ifdef __SPIR__ +#if defined(__SPIR__) || defined(__SPIRV__) #include "device.h" #include diff --git a/libdevice/sanitizer_utils.cpp b/libdevice/sanitizer_utils.cpp index f668d7b65a168..c560d1d731bdc 100644 --- a/libdevice/sanitizer_utils.cpp +++ b/libdevice/sanitizer_utils.cpp @@ -34,7 +34,7 @@ DeviceGlobal __DeviceSanitizerReportMem; DeviceGlobal __DeviceType; -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) #ifdef __SYCL_DEVICE_ONLY__ #define __USE_SPIR_BUILTIN__ 1 diff --git a/libdevice/spirv_vars.h b/libdevice/spirv_vars.h index 0387c322965ce..0515fa9b15f34 100644 --- a/libdevice/spirv_vars.h +++ b/libdevice/spirv_vars.h @@ -11,7 +11,7 @@ #include "device.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) #include #include @@ -52,10 +52,10 @@ DEVICE_EXTERNAL inline size_t __spirv_LocalInvocationId_z() { return __spirv_BuiltInLocalInvocationId.z; } -#ifndef __SPIR__ +#if !defined(__SPIR__) && !defined(__SPIRV__) const size_t_vec __spirv_BuiltInGlobalInvocationId{}; const size_t_vec __spirv_BuiltInLocalInvocationId{}; -#endif // __SPIR__ +#endif // !__SPIR__ && !__SPIRV__ -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ #endif // __LIBDEVICE_SPIRV_VARS_H diff --git a/libdevice/wrapper.h b/libdevice/wrapper.h index c3ec6ec1fa785..bbc0cd4f1ca87 100644 --- a/libdevice/wrapper.h +++ b/libdevice/wrapper.h @@ -11,7 +11,7 @@ #include "device.h" -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) #include #include @@ -29,5 +29,5 @@ void __devicelib_assert_fail(const char *expr, const char *file, int32_t line, const char *func, uint64_t gid0, uint64_t gid1, uint64_t gid2, uint64_t lid0, uint64_t lid1, uint64_t lid2); -#endif // __SPIR__ || __NVPTX__ +#endif // __SPIR__ || __SPIRV__ || __NVPTX__ #endif // __LIBDEVICE_WRAPPER_H__ diff --git a/sycl/include/sycl/ext/intel/math.hpp b/sycl/include/sycl/ext/intel/math.hpp index 1b3be573d0a16..63d25e51064f3 100644 --- a/sycl/include/sycl/ext/intel/math.hpp +++ b/sycl/include/sycl/ext/intel/math.hpp @@ -13,7 +13,7 @@ // _iml_half_internal is internal representation for fp16 type used in intel // math device library. The definition here should align with definition in // https://github.com/intel/llvm/blob/sycl/libdevice/imf_half.hpp -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) using _iml_half_internal = _Float16; #else using _iml_half_internal = uint16_t; diff --git a/sycl/include/sycl/ext/oneapi/bfloat16.hpp b/sycl/include/sycl/ext/oneapi/bfloat16.hpp index 3a0e1259a1ff8..3a16dcd244b4c 100644 --- a/sycl/include/sycl/ext/oneapi/bfloat16.hpp +++ b/sycl/include/sycl/ext/oneapi/bfloat16.hpp @@ -102,7 +102,7 @@ class bfloat16 { } static float to_float(const detail::Bfloat16StorageT &a) { -#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) +#if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__)) return __devicelib_ConvertBF16ToFINTEL(a); #else union { @@ -155,7 +155,7 @@ class bfloat16 { detail::Bfloat16StorageT res; asm("neg.bf16 %0, %1;" : "=h"(res) : "h"(lhs.value)); return detail::bitsToBfloat16(res); -#elif defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) +#elif defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__)) return bfloat16{-__devicelib_ConvertBF16ToFINTEL(lhs.value)}; #else return bfloat16{-to_float(lhs.value)}; diff --git a/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp index 079a637580b93..e41cc9969c71a 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp @@ -146,7 +146,7 @@ inline std::enable_if_t> && get_ballot_group(Group group, bool predicate) { (void)group; #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__SPIR__) || defined(__NVPTX__) +#if defined(__SPIR__) || defined(__SPIRV__) || defined(__NVPTX__) // ballot_group partitions into two groups using the predicate // Membership mask for one group is negation of the other sub_group_mask mask = sycl::ext::oneapi::group_ballot(group, predicate); diff --git a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp index f46984334f6f4..73959795961fb 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp @@ -77,11 +77,12 @@ namespace ext::oneapi::experimental { // template int printf(const FormatT *__format, Args... args) { -#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) +#if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__)) return __spirv_ocl_printf(__format, args...); #else return ::printf(__format, args...); -#endif // defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) +#endif // defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || + // defined(__SPIRV__)) } namespace native { diff --git a/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp index b774225049420..d0433d9add01c 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp @@ -140,7 +140,7 @@ namespace this_kernel { inline opportunistic_group get_opportunistic_group() { #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) // TODO: It may be wiser to call the intrinsic than rely on this_group() sycl::sub_group sg = sycl::ext::oneapi::experimental::this_sub_group(); sub_group_mask mask = sycl::ext::oneapi::group_ballot(sg, true); diff --git a/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp index 22292a95f76fd..ebc6243248deb 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp @@ -142,7 +142,7 @@ inline std::enable_if_t> && get_tangle_group(Group group) { (void)group; #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) // All SPIR-V devices that we currently target execute in SIMD fashion, // and so the group of work-items in converged control flow is implicit. // We store the mask here because it is required to calculate IDs, not diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp index f0d8b4d19a412..d130600c7061a 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp @@ -50,7 +50,7 @@ struct joint_matrix { #elif defined(__HIP_PLATFORM_AMD_MFMA__) sycl::ext::oneapi::detail::joint_matrix_hip matrix_impl; -#elif defined(__SPIR__) +#elif defined(__SPIR__) || defined(__SPIRV__) __spv::__spirv_JointMatrixINTEL< T, Rows, Cols, spv_matrix_layout_traits::value, spv_scope_traits::value, spv_matrix_use_traits::value> *spvm; @@ -74,10 +74,10 @@ struct joint_matrix { #endif } #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) joint_matrix(const joint_matrix &other) = delete; joint_matrix &operator=(const joint_matrix &rhs) = delete; -#endif // defined(__SPIR__) +#endif // defined(__SPIR__) || defined(__SPIRV__) #endif }; diff --git a/sycl/test-e2e/DeviceLib/built-ins/printf.cpp b/sycl/test-e2e/DeviceLib/built-ins/printf.cpp index 718d9b292044b..8ad3a6b1a670d 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/printf.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/printf.cpp @@ -56,7 +56,7 @@ int main() { // Vectors sycl::vec v4{5, 6, 7, 8}; -#if defined(__SYCL_DEVICE_ONLY__) && defined(__SPIR__) +#if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__)) // On SPIRV devices, vectors can be printed via native OpenCL types: using ocl_int4 = sycl::vec::vector_t; { diff --git a/sycl/test-e2e/DeviceLib/imf_utils.hpp b/sycl/test-e2e/DeviceLib/imf_utils.hpp index 846011e77e24a..8559050196413 100644 --- a/sycl/test-e2e/DeviceLib/imf_utils.hpp +++ b/sycl/test-e2e/DeviceLib/imf_utils.hpp @@ -7,7 +7,7 @@ #include #include -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) typedef _Float16 _iml_half_internal; #else typedef uint16_t _iml_half_internal; @@ -213,7 +213,7 @@ void test3(sycl::queue &q, std::initializer_list Input1, #define F3(Name) [](auto x, auto y, auto z) { return (Name)(x, y, z); } #define F3T(T, Name) \ [](auto x, auto y, auto z) { return __builtin_bit_cast(T, (Name)(x, y, z)); } -#if defined(__SPIR__) +#if defined(__SPIR__) || defined(__SPIRV__) #define F_Half1(Name) \ [](uint16_t x) { return (Name)(__builtin_bit_cast(_Float16, x)); } #define F_Half2(Name) \