diff options
Diffstat (limited to 'intern/cycles/util/util_atomic.h')
| -rw-r--r-- | intern/cycles/util/util_atomic.h | 32 |
1 files changed, 30 insertions, 2 deletions
diff --git a/intern/cycles/util/util_atomic.h b/intern/cycles/util/util_atomic.h index f3c7ae546a0..477b667a6fe 100644 --- a/intern/cycles/util/util_atomic.h +++ b/intern/cycles/util/util_atomic.h @@ -23,12 +23,13 @@ #include "atomic_ops.h" #define atomic_add_and_fetch_float(p, x) atomic_add_and_fetch_fl((p), (x)) +#define atomic_compare_and_swap_float(p, old_val, new_val) atomic_cas_float((p), (old_val), (new_val)) #define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1) #define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_add_uint32((p), -1) #define CCL_LOCAL_MEM_FENCE 0 -#define ccl_barrier(flags) (void)0 +#define ccl_barrier(flags) ((void) 0) #else /* __KERNEL_GPU__ */ @@ -57,6 +58,20 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so return new_value.float_value; } +ccl_device_inline float atomic_compare_and_swap_float(volatile ccl_global float *dest, + const float old_val, const float new_val) +{ + union { + unsigned int int_value; + float float_value; + } new_value, prev_value, result; + prev_value.float_value = old_val; + new_value.float_value = new_val; + result.int_value = atomic_cmpxchg((volatile ccl_global unsigned int *)dest, + prev_value.int_value, new_value.int_value); + return result.float_value; +} + #define atomic_fetch_and_add_uint32(p, x) atomic_add((p), (x)) #define atomic_fetch_and_inc_uint32(p) atomic_inc((p)) #define atomic_fetch_and_dec_uint32(p) atomic_dec((p)) @@ -75,6 +90,19 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so #define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1) #define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_sub_uint32((p), 1) +ccl_device_inline float atomic_compare_and_swap_float(volatile float *dest, + const float old_val, const float new_val) +{ + union { + unsigned int int_value; + float float_value; + } new_value, prev_value, result; + prev_value.float_value = old_val; + new_value.float_value = new_val; + result.int_value = atomicCAS((unsigned int *)dest, prev_value.int_value,new_value.int_value); + return result.float_value; +} + #define CCL_LOCAL_MEM_FENCE #define ccl_barrier(flags) __syncthreads() @@ -82,4 +110,4 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so #endif /* __KERNEL_GPU__ */ -#endif /* __UTIL_ATOMIC_H__ */ +#endif /* __UTIL_ATOMIC_H__ */ |