diff options
Diffstat (limited to 'intern/cycles/util/atomic.h')
| -rw-r--r-- | intern/cycles/util/atomic.h | 110 |
1 files changed, 110 insertions, 0 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__ */ |