17 #ifndef __KERNEL_QUEUE_H__
18 #define __KERNEL_QUEUE_H__
25 #ifdef __KERNEL_OPENCL__
26 # pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
27 # pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
42 (queue_number * queue_size);
43 queues[my_queue_index] = ray_index;
61 int ray_index = queues[queue_number * queuesize + thread_index];
92 (
ccl_global uint *)&Queue_index[queue_number], *local_queue_atomics);
98 unsigned int my_gqidx = queue_number * queuesize + (*local_queue_atomics) + lqidx;
99 Queue_data[my_gqidx] = ray_index;
117 (
ccl_global uint *)&global_queue_atomics[queue_number], local_queue_atomics[queue_number]);
127 int my_gqidx = queuesize * queue_number + lqidx + global_per_queue_offset[queue_number];
142 return queues[index + queue_number * queue_size];
ATOMIC_INLINE uint32_t atomic_fetch_and_add_uint32(uint32_t *p, uint32_t x)
ccl_device_inline uint ccl_local_id(uint d)
ccl_device_inline uint ccl_local_size(uint d)
#define CCL_NAMESPACE_END
ccl_device unsigned int get_global_queue_index(int queue_number, int queuesize, unsigned int lqidx, ccl_local_param unsigned int *global_per_queue_offset)
ccl_device unsigned int get_local_queue_index(int queue_number, ccl_local_param unsigned int *local_queue_atomics)
CCL_NAMESPACE_BEGIN ccl_device void enqueue_ray_index(int ray_index, int queue_number, ccl_global int *queues, int queue_size, ccl_global int *queue_index)
ccl_device int dequeue_ray_index(int queue_number, ccl_global int *queues, int queue_size, ccl_global int *queue_index)
ccl_device int get_ray_index(KernelGlobals *kg, int thread_index, int queue_number, ccl_global int *queues, int queuesize, int empty_queue)
ccl_device unsigned int get_global_per_queue_offset(int queue_number, ccl_local_param unsigned int *local_queue_atomics, ccl_global int *global_queue_atomics)
ccl_device void enqueue_ray_index_local(int ray_index, int queue_number, char enqueue_flag, int queuesize, ccl_local_param unsigned int *local_queue_atomics, ccl_global int *Queue_data, ccl_global int *Queue_index)
__kernel void ccl_constant KernelData ccl_global void ccl_global char ccl_global int * queue_index
#define CCL_LOCAL_MEM_FENCE
#define atomic_fetch_and_dec_uint32(p)
#define atomic_fetch_and_inc_uint32(p)
#define ccl_barrier(flags)