diff options
Diffstat (limited to 'intern/cycles/util')
30 files changed, 185 insertions, 59 deletions
diff --git a/intern/cycles/util/atomic.h b/intern/cycles/util/atomic.h index f89eb28b0b7..1ebf085ae13 100644 --- a/intern/cycles/util/atomic.h +++ b/intern/cycles/util/atomic.h @@ -106,6 +106,116 @@ ccl_device_inline float atomic_compare_and_swap_float(volatile ccl_global float # endif /* __KERNEL_METAL__ */ +# ifdef __KERNEL_ONEAPI__ + +ccl_device_inline float atomic_add_and_fetch_float(ccl_global float *p, float x) +{ + sycl::atomic_ref<float, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*p); + return atomic.fetch_add(x); +} + +ccl_device_inline float atomic_compare_and_swap_float(ccl_global float *source, + float old_val, + float new_val) +{ + sycl::atomic_ref<float, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*source); + atomic.compare_exchange_weak(old_val, new_val); + return old_val; +} + +ccl_device_inline unsigned int atomic_fetch_and_add_uint32(ccl_global unsigned int *p, + unsigned int x) +{ + sycl::atomic_ref<unsigned int, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*p); + return atomic.fetch_add(x); +} + +ccl_device_inline int atomic_fetch_and_add_uint32(ccl_global int *p, int x) +{ + sycl::atomic_ref<int, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*p); + return atomic.fetch_add(x); +} + +ccl_device_inline unsigned int atomic_fetch_and_sub_uint32(ccl_global unsigned int *p, + unsigned int x) +{ + sycl::atomic_ref<unsigned int, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*p); + return atomic.fetch_sub(x); +} + +ccl_device_inline int atomic_fetch_and_sub_uint32(ccl_global int *p, int x) +{ + sycl::atomic_ref<int, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*p); + return atomic.fetch_sub(x); +} + +ccl_device_inline unsigned int atomic_fetch_and_inc_uint32(ccl_global unsigned int *p) +{ + return atomic_fetch_and_add_uint32(p, 1); +} + +ccl_device_inline int atomic_fetch_and_inc_uint32(ccl_global int *p) +{ + return atomic_fetch_and_add_uint32(p, 1); +} + +ccl_device_inline unsigned int atomic_fetch_and_dec_uint32(ccl_global unsigned int *p) +{ + return atomic_fetch_and_sub_uint32(p, 1); +} + +ccl_device_inline int atomic_fetch_and_dec_uint32(ccl_global int *p) +{ + return atomic_fetch_and_sub_uint32(p, 1); +} + +ccl_device_inline unsigned int atomic_fetch_and_or_uint32(ccl_global unsigned int *p, + unsigned int x) +{ + sycl::atomic_ref<unsigned int, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*p); + return atomic.fetch_or(x); +} + +ccl_device_inline int atomic_fetch_and_or_uint32(ccl_global int *p, int x) +{ + sycl::atomic_ref<int, + sycl::memory_order::relaxed, + sycl::memory_scope::device, + sycl::access::address_space::ext_intel_global_device_space> + atomic(*p); + return atomic.fetch_or(x); +} + +# endif /* __KERNEL_ONEAPI__ */ + #endif /* __KERNEL_GPU__ */ #endif /* __UTIL_ATOMIC_H__ */ diff --git a/intern/cycles/util/half.h b/intern/cycles/util/half.h index 434bc12d670..c668638eb02 100644 --- a/intern/cycles/util/half.h +++ b/intern/cycles/util/half.h @@ -35,7 +35,7 @@ ccl_device_inline float half_to_float(half h_in) #else /* CUDA has its own half data type, no need to define then */ -# if !defined(__KERNEL_CUDA__) && !defined(__KERNEL_HIP__) +# if !defined(__KERNEL_CUDA__) && !defined(__KERNEL_HIP__) && !defined(__KERNEL_ONEAPI__) /* Implementing this as a class rather than a typedef so that the compiler can tell it apart from * unsigned shorts. */ class half { @@ -73,7 +73,7 @@ struct half4 { ccl_device_inline half float_to_half_image(float f) { -#if defined(__KERNEL_METAL__) +#if defined(__KERNEL_METAL__) || defined(__KERNEL_ONEAPI__) return half(min(f, 65504.0f)); #elif defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__) return __float2half(min(f, 65504.0f)); @@ -103,6 +103,8 @@ ccl_device_inline float half_to_float_image(half h) { #if defined(__KERNEL_METAL__) return half_to_float(h); +#elif defined(__KERNEL_ONEAPI__) + return float(h); #elif defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__) return __half2float(h); #else @@ -136,7 +138,7 @@ ccl_device_inline float4 half4_to_float4_image(const half4 h) ccl_device_inline half float_to_half_display(const float f) { -#if defined(__KERNEL_METAL__) +#if defined(__KERNEL_METAL__) || defined(__KERNEL_ONEAPI__) return half(min(f, 65504.0f)); #elif defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__) return __float2half(min(f, 65504.0f)); diff --git a/intern/cycles/util/math.h b/intern/cycles/util/math.h index f1f627588c5..af2f1ea092d 100644 --- a/intern/cycles/util/math.h +++ b/intern/cycles/util/math.h @@ -79,7 +79,7 @@ CCL_NAMESPACE_BEGIN /* Scalar */ -#ifndef __HIP__ +#if !defined(__HIP__) && !defined(__KERNEL_ONEAPI__) # ifdef _WIN32 ccl_device_inline float fmaxf(float a, float b) { @@ -92,12 +92,18 @@ ccl_device_inline float fminf(float a, float b) } # endif /* _WIN32 */ -#endif /* __HIP__ */ +#endif /* __HIP__, __KERNEL_ONEAPI__ */ -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) +# ifndef __KERNEL_ONEAPI__ using std::isfinite; using std::isnan; using std::sqrt; +# else +using sycl::sqrt; +# define isfinite(x) sycl::isfinite((x)) +# define isnan(x) sycl::isnan((x)) +# endif ccl_device_inline int abs(int x) { @@ -793,6 +799,8 @@ ccl_device_inline uint popcount(uint x) return i & 1; } # endif +#elif defined(__KERNEL_ONEAPI__) +# define popcount(x) sycl::popcount(x) #elif defined(__KERNEL_HIP__) /* Use popcll to support 64-bit wave for pre-RDNA AMD GPUs */ # define popcount(x) __popcll(x) @@ -806,6 +814,8 @@ ccl_device_inline uint count_leading_zeros(uint x) return __clz(x); #elif defined(__KERNEL_METAL__) return clz(x); +#elif defined(__KERNEL_ONEAPI__) + return sycl::clz(x); #else assert(x != 0); # ifdef _MSC_VER @@ -824,6 +834,8 @@ ccl_device_inline uint count_trailing_zeros(uint x) return (__ffs(x) - 1); #elif defined(__KERNEL_METAL__) return ctz(x); +#elif defined(__KERNEL_ONEAPI__) + return sycl::ctz(x); #else assert(x != 0); # ifdef _MSC_VER diff --git a/intern/cycles/util/types_float2.h b/intern/cycles/util/types_float2.h index d8b2efb7b4b..07b9ec0986b 100644 --- a/intern/cycles/util/types_float2.h +++ b/intern/cycles/util/types_float2.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct float2 { float x, y; @@ -20,7 +20,7 @@ struct float2 { ccl_device_inline float2 make_float2(float x, float y); ccl_device_inline void print_float2(const char *label, const float2 &a); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_float2_impl.h b/intern/cycles/util/types_float2_impl.h index d67ec946b79..45fc90c52bd 100644 --- a/intern/cycles/util/types_float2_impl.h +++ b/intern/cycles/util/types_float2_impl.h @@ -14,7 +14,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) __forceinline float float2::operator[](int i) const { util_assert(i >= 0); @@ -39,7 +39,7 @@ ccl_device_inline void print_float2(const char *label, const float2 &a) { printf("%s: %.8f %.8f\n", label, (double)a.x, (double)a.y); } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_float3.h b/intern/cycles/util/types_float3.h index 060c2ac4152..c7900acaa69 100644 --- a/intern/cycles/util/types_float3.h +++ b/intern/cycles/util/types_float3.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) struct ccl_try_align(16) float3 { # ifdef __KERNEL_SSE__ @@ -40,7 +40,7 @@ struct ccl_try_align(16) float3 ccl_device_inline float3 make_float3(float f); ccl_device_inline float3 make_float3(float x, float y, float z); ccl_device_inline void print_float3(const char *label, const float3 &a); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) */ /* Smaller float3 for storage. For math operations this must be converted to float3, so that on the * CPU SIMD instructions can be used. */ diff --git a/intern/cycles/util/types_float3_impl.h b/intern/cycles/util/types_float3_impl.h index f5ffc48c1be..2e6e864c8ea 100644 --- a/intern/cycles/util/types_float3_impl.h +++ b/intern/cycles/util/types_float3_impl.h @@ -14,7 +14,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) # ifdef __KERNEL_SSE__ __forceinline float3::float3() { @@ -83,7 +83,7 @@ ccl_device_inline void print_float3(const char *label, const float3 &a) { printf("%s: %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z); } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_float4.h b/intern/cycles/util/types_float4.h index 68ba787dac0..27453bf39e4 100644 --- a/intern/cycles/util/types_float4.h +++ b/intern/cycles/util/types_float4.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct int4; struct ccl_try_align(16) float4 @@ -43,7 +43,7 @@ ccl_device_inline float4 make_float4(float f); ccl_device_inline float4 make_float4(float x, float y, float z, float w); ccl_device_inline float4 make_float4(const int4 &i); ccl_device_inline void print_float4(const char *label, const float4 &a); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_float4_impl.h b/intern/cycles/util/types_float4_impl.h index de2e7cb7061..d7858f744e3 100644 --- a/intern/cycles/util/types_float4_impl.h +++ b/intern/cycles/util/types_float4_impl.h @@ -14,7 +14,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) # ifdef __KERNEL_SSE__ __forceinline float4::float4() { @@ -89,7 +89,7 @@ ccl_device_inline void print_float4(const char *label, const float4 &a) { printf("%s: %.8f %.8f %.8f %.8f\n", label, (double)a.x, (double)a.y, (double)a.z, (double)a.w); } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_float8.h b/intern/cycles/util/types_float8.h index 99f9ec9b867..d71149946f7 100644 --- a/intern/cycles/util/types_float8.h +++ b/intern/cycles/util/types_float8.h @@ -11,7 +11,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct ccl_try_align(32) float8 { @@ -43,7 +43,7 @@ struct ccl_try_align(32) float8 ccl_device_inline float8 make_float8(float f); ccl_device_inline float8 make_float8(float a, float b, float c, float d, float e, float f, float g, float h); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_float8_impl.h b/intern/cycles/util/types_float8_impl.h index 19818976b50..0694f5205a5 100644 --- a/intern/cycles/util/types_float8_impl.h +++ b/intern/cycles/util/types_float8_impl.h @@ -15,7 +15,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) # ifdef __KERNEL_AVX2__ __forceinline float8::float8() { @@ -81,7 +81,7 @@ make_float8(float a, float b, float c, float d, float e, float f, float g, float return r; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_int2.h b/intern/cycles/util/types_int2.h index 4daf387d9cf..bf69cddc653 100644 --- a/intern/cycles/util/types_int2.h +++ b/intern/cycles/util/types_int2.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct int2 { int x, y; @@ -19,7 +19,7 @@ struct int2 { }; ccl_device_inline int2 make_int2(int x, int y); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_int2_impl.h b/intern/cycles/util/types_int2_impl.h index 7989c4d5506..7bdc77369ee 100644 --- a/intern/cycles/util/types_int2_impl.h +++ b/intern/cycles/util/types_int2_impl.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) int int2::operator[](int i) const { util_assert(i >= 0); @@ -30,7 +30,7 @@ ccl_device_inline int2 make_int2(int x, int y) int2 a = {x, y}; return a; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_int3.h b/intern/cycles/util/types_int3.h index ad9bcb39bbe..f88ff22ac35 100644 --- a/intern/cycles/util/types_int3.h +++ b/intern/cycles/util/types_int3.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct ccl_try_align(16) int3 { # ifdef __KERNEL_SSE__ @@ -40,7 +40,7 @@ struct ccl_try_align(16) int3 ccl_device_inline int3 make_int3(int i); ccl_device_inline int3 make_int3(int x, int y, int z); ccl_device_inline void print_int3(const char *label, const int3 &a); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_int3_impl.h b/intern/cycles/util/types_int3_impl.h index 4cfc1cf2987..1c49e97ad32 100644 --- a/intern/cycles/util/types_int3_impl.h +++ b/intern/cycles/util/types_int3_impl.h @@ -14,7 +14,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) # ifdef __KERNEL_SSE__ __forceinline int3::int3() { @@ -84,7 +84,7 @@ ccl_device_inline void print_int3(const char *label, const int3 &a) { printf("%s: %d %d %d\n", label, a.x, a.y, a.z); } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_int4.h b/intern/cycles/util/types_int4.h index f35632fb52f..9d557c01344 100644 --- a/intern/cycles/util/types_int4.h +++ b/intern/cycles/util/types_int4.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct float3; struct float4; @@ -46,7 +46,7 @@ ccl_device_inline int4 make_int4(int x, int y, int z, int w); ccl_device_inline int4 make_int4(const float3 &f); ccl_device_inline int4 make_int4(const float4 &f); ccl_device_inline void print_int4(const char *label, const int4 &a); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_int4_impl.h b/intern/cycles/util/types_int4_impl.h index adb4a4cebac..11e1ede6705 100644 --- a/intern/cycles/util/types_int4_impl.h +++ b/intern/cycles/util/types_int4_impl.h @@ -14,7 +14,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) # ifdef __KERNEL_SSE__ __forceinline int4::int4() { @@ -83,6 +83,8 @@ ccl_device_inline int4 make_int4(const float3 &f) { # ifdef __KERNEL_SSE__ int4 a(_mm_cvtps_epi32(f.m128)); +# elif defined(__KERNEL_ONEAPI__) + int4 a = {(int)f.x, (int)f.y, (int)f.z, 0}; # else int4 a = {(int)f.x, (int)f.y, (int)f.z, (int)f.w}; # endif @@ -103,7 +105,7 @@ ccl_device_inline void print_int4(const char *label, const int4 &a) { printf("%s: %d %d %d %d\n", label, a.x, a.y, a.z, a.w); } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uchar2.h b/intern/cycles/util/types_uchar2.h index 445fa8dd703..0b3c9bd0331 100644 --- a/intern/cycles/util/types_uchar2.h +++ b/intern/cycles/util/types_uchar2.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct uchar2 { uchar x, y; @@ -19,7 +19,7 @@ struct uchar2 { }; ccl_device_inline uchar2 make_uchar2(uchar x, uchar y); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uchar2_impl.h b/intern/cycles/util/types_uchar2_impl.h index cec1c679050..a7254d5eaf2 100644 --- a/intern/cycles/util/types_uchar2_impl.h +++ b/intern/cycles/util/types_uchar2_impl.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) uchar uchar2::operator[](int i) const { util_assert(i >= 0); @@ -30,7 +30,7 @@ ccl_device_inline uchar2 make_uchar2(uchar x, uchar y) uchar2 a = {x, y}; return a; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uchar3.h b/intern/cycles/util/types_uchar3.h index 1ebd86441c3..fc213502ada 100644 --- a/intern/cycles/util/types_uchar3.h +++ b/intern/cycles/util/types_uchar3.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct uchar3 { uchar x, y, z; @@ -19,7 +19,7 @@ struct uchar3 { }; ccl_device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uchar3_impl.h b/intern/cycles/util/types_uchar3_impl.h index 0656baa3da4..0c24ffb488a 100644 --- a/intern/cycles/util/types_uchar3_impl.h +++ b/intern/cycles/util/types_uchar3_impl.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) uchar uchar3::operator[](int i) const { util_assert(i >= 0); @@ -30,7 +30,7 @@ ccl_device_inline uchar3 make_uchar3(uchar x, uchar y, uchar z) uchar3 a = {x, y, z}; return a; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uchar4.h b/intern/cycles/util/types_uchar4.h index 2ac4fb56cbb..a2a2c945aaa 100644 --- a/intern/cycles/util/types_uchar4.h +++ b/intern/cycles/util/types_uchar4.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct uchar4 { uchar x, y, z, w; @@ -19,7 +19,7 @@ struct uchar4 { }; ccl_device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uchar4_impl.h b/intern/cycles/util/types_uchar4_impl.h index b3e8abfe873..8ec6213a37d 100644 --- a/intern/cycles/util/types_uchar4_impl.h +++ b/intern/cycles/util/types_uchar4_impl.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) uchar uchar4::operator[](int i) const { util_assert(i >= 0); @@ -30,7 +30,7 @@ ccl_device_inline uchar4 make_uchar4(uchar x, uchar y, uchar z, uchar w) uchar4 a = {x, y, z, w}; return a; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uint2.h b/intern/cycles/util/types_uint2.h index e3254b9f0e1..faa0955f903 100644 --- a/intern/cycles/util/types_uint2.h +++ b/intern/cycles/util/types_uint2.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct uint2 { uint x, y; @@ -19,7 +19,7 @@ struct uint2 { }; ccl_device_inline uint2 make_uint2(uint x, uint y); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uint2_impl.h b/intern/cycles/util/types_uint2_impl.h index e67134a011e..cac0ba6b531 100644 --- a/intern/cycles/util/types_uint2_impl.h +++ b/intern/cycles/util/types_uint2_impl.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) __forceinline uint uint2::operator[](uint i) const { util_assert(i < 2); @@ -28,7 +28,7 @@ ccl_device_inline uint2 make_uint2(uint x, uint y) uint2 a = {x, y}; return a; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uint3.h b/intern/cycles/util/types_uint3.h index 885a8fb84ce..3ff87bfc791 100644 --- a/intern/cycles/util/types_uint3.h +++ b/intern/cycles/util/types_uint3.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct uint3 { uint x, y, z; @@ -19,7 +19,7 @@ struct uint3 { }; ccl_device_inline uint3 make_uint3(uint x, uint y, uint z); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uint3_impl.h b/intern/cycles/util/types_uint3_impl.h index f4d3d72469c..221883a1adb 100644 --- a/intern/cycles/util/types_uint3_impl.h +++ b/intern/cycles/util/types_uint3_impl.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) __forceinline uint uint3::operator[](uint i) const { util_assert(i < 3); @@ -28,7 +28,7 @@ ccl_device_inline uint3 make_uint3(uint x, uint y, uint z) uint3 a = {x, y, z}; return a; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uint4.h b/intern/cycles/util/types_uint4.h index d582b91d2a0..504095b2383 100644 --- a/intern/cycles/util/types_uint4.h +++ b/intern/cycles/util/types_uint4.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct uint4 { uint x, y, z, w; @@ -19,7 +19,7 @@ struct uint4 { }; ccl_device_inline uint4 make_uint4(uint x, uint y, uint z, uint w); -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_uint4_impl.h b/intern/cycles/util/types_uint4_impl.h index 98a4c5e9fe9..d78db944a1f 100644 --- a/intern/cycles/util/types_uint4_impl.h +++ b/intern/cycles/util/types_uint4_impl.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) __forceinline uint uint4::operator[](uint i) const { util_assert(i < 3); @@ -28,7 +28,7 @@ ccl_device_inline uint4 make_uint4(uint x, uint y, uint z, uint w) uint4 a = {x, y, z, w}; return a; } -#endif /* __KERNEL_GPU__ */ +#endif /* !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) */ CCL_NAMESPACE_END diff --git a/intern/cycles/util/types_ushort4.h b/intern/cycles/util/types_ushort4.h index 1766c6bf734..9a6e12095ba 100644 --- a/intern/cycles/util/types_ushort4.h +++ b/intern/cycles/util/types_ushort4.h @@ -10,7 +10,7 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_GPU__ +#if !defined(__KERNEL_GPU__) || defined(__KERNEL_ONEAPI__) struct ushort4 { uint16_t x, y, z, w; |