diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-07-03 11:56:40 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-07-03 12:03:56 +0300 |
commit | b9f89b164785f28b023af071d696c298deba4f63 (patch) | |
tree | 4fd045caacf2850d34820ac4d07422ca83cc2fba /intern/cycles/kernel/kernel_queues.h | |
parent | 80f344fd952a42fd7250e1eca67c8f2c19d8897b (diff) |
Cycles: Code cleanup in split kernel, whitespaces
Diffstat (limited to 'intern/cycles/kernel/kernel_queues.h')
-rw-r--r-- | intern/cycles/kernel/kernel_queues.h | 79 |
1 files changed, 36 insertions, 43 deletions
diff --git a/intern/cycles/kernel/kernel_queues.h b/intern/cycles/kernel/kernel_queues.h index 9e65e2b0768..cf5614b8a86 100644 --- a/intern/cycles/kernel/kernel_queues.h +++ b/intern/cycles/kernel/kernel_queues.h @@ -27,15 +27,14 @@ /* * Enqueue ray index into the queue */ -ccl_device void enqueue_ray_index ( - int ray_index, /* Ray index to be enqueued */ - int queue_number, /* Queue in which the ray index should be enqueued*/ - ccl_global int *queues, /* Buffer of all queues */ - int queue_size, /* Size of each queue */ - ccl_global int *queue_index /* Array of size num_queues; Used for atomic increment */ - ) +ccl_device void enqueue_ray_index( + int ray_index, /* Ray index to be enqueued. */ + int queue_number, /* Queue in which the ray index should be enqueued. */ + ccl_global int *queues, /* Buffer of all queues. */ + int queue_size, /* Size of each queue. */ + ccl_global int *queue_index) /* Array of size num_queues; Used for atomic increment. */ { - /* This thread's queue index */ + /* This thread's queue index. */ int my_queue_index = atomic_inc(&queue_index[queue_number]) + (queue_number * queue_size); queues[my_queue_index] = ray_index; } @@ -47,52 +46,48 @@ ccl_device void enqueue_ray_index ( * i.e All ray's in the queue has been successfully allocated and there * is no more ray to allocate to other threads. */ -ccl_device int get_ray_index ( - int thread_index, /* Global thread index */ - int queue_number, /* Queue to operate on */ - ccl_global int *queues, /* Buffer of all queues */ - int queuesize, /* Size of a queue */ - int empty_queue /* Empty the queue slot as soon as we fetch the ray index */ - ) +ccl_device int get_ray_index( + int thread_index, /* Global thread index. */ + int queue_number, /* Queue to operate on. */ + ccl_global int *queues, /* Buffer of all queues. */ + int queuesize, /* Size of a queue. */ + int empty_queue) /* Empty the queue slot as soon as we fetch the ray index. */ { int ray_index = queues[queue_number * queuesize + thread_index]; - if(empty_queue && ray_index != QUEUE_EMPTY_SLOT) { queues[queue_number * queuesize + thread_index] = QUEUE_EMPTY_SLOT; } - return ray_index; } -/* The following functions are to realize Local memory variant of enqueue ray index function */ +/* The following functions are to realize Local memory variant of enqueue ray index function. */ -/* All threads should call this function */ +/* All threads should call this function. */ ccl_device void enqueue_ray_index_local( - int ray_index, /* Ray index to enqueue*/ - int queue_number, /* Queue in which to enqueue ray index */ - char enqueue_flag, /* True for threads whose ray index has to be enqueued */ - int queuesize, /* queue size */ - ccl_local unsigned int *local_queue_atomics, /* To to local queue atomics */ - ccl_global int *Queue_data, /* Queues */ - ccl_global int *Queue_index /* To do global queue atomics */ - ) + int ray_index, /* Ray index to enqueue. */ + int queue_number, /* Queue in which to enqueue ray index. */ + char enqueue_flag, /* True for threads whose ray index has to be enqueued. */ + int queuesize, /* queue size. */ + ccl_local unsigned int *local_queue_atomics, /* To to local queue atomics. */ + ccl_global int *Queue_data, /* Queues. */ + ccl_global int *Queue_index) /* To do global queue atomics. */ { int lidx = get_local_id(1) * get_local_size(0) + get_local_id(0); - /* Get local queue id */ + /* Get local queue id .*/ unsigned int lqidx; if(enqueue_flag) { lqidx = atomic_inc(local_queue_atomics); } barrier(CLK_LOCAL_MEM_FENCE); - /* Get global queue offset */ + /* Get global queue offset. */ if(lidx == 0) { *local_queue_atomics = atomic_add(&Queue_index[queue_number], *local_queue_atomics); } barrier(CLK_LOCAL_MEM_FENCE); - /* Get global queue index and enqueue ray */ + /* Get global queue index and enqueue ray. */ if(enqueue_flag) { unsigned int my_gqidx = queue_number * queuesize + (*local_queue_atomics) + lqidx; Queue_data[my_gqidx] = ray_index; @@ -100,30 +95,28 @@ ccl_device void enqueue_ray_index_local( } ccl_device unsigned int get_local_queue_index( - int queue_number, /* Queue in which to enqueue the ray; -1 if no queue */ - ccl_local unsigned int *local_queue_atomics - ) + int queue_number, /* Queue in which to enqueue the ray; -1 if no queue */ + ccl_local unsigned int *local_queue_atomics) { int my_lqidx = atomic_inc(&local_queue_atomics[queue_number]); return my_lqidx; } ccl_device unsigned int get_global_per_queue_offset( - int queue_number, - ccl_local unsigned int *local_queue_atomics, - ccl_global int* global_queue_atomics - ) + int queue_number, + ccl_local unsigned int *local_queue_atomics, + ccl_global int* global_queue_atomics) { - unsigned int queue_offset = atomic_add((&global_queue_atomics[queue_number]), local_queue_atomics[queue_number]); + unsigned int queue_offset = atomic_add(&global_queue_atomics[queue_number], + local_queue_atomics[queue_number]); return queue_offset; } ccl_device unsigned int get_global_queue_index( - int queue_number, - int queuesize, - unsigned int lqidx, - ccl_local unsigned int * global_per_queue_offset - ) + int queue_number, + int queuesize, + unsigned int lqidx, + ccl_local unsigned int * global_per_queue_offset) { int my_gqidx = queuesize * queue_number + lqidx + global_per_queue_offset[queue_number]; return my_gqidx; |