From b9f89b164785f28b023af071d696c298deba4f63 Mon Sep 17 00:00:00 2001 From: Sergey Sharybin Date: Fri, 3 Jul 2015 10:56:40 +0200 Subject: [PATCH] Cycles: Code cleanup in split kernel, whitespaces --- intern/cycles/kernel/kernel_queues.h | 79 +++++++++---------- .../split/kernel_background_buffer_update.h | 13 ++- intern/cycles/kernel/split/kernel_data_init.h | 11 +-- .../split/kernel_next_iteration_setup.h | 3 +- .../kernel/split/kernel_scene_intersect.h | 3 +- .../kernel/split/kernel_shadow_blocked.h | 2 +- 6 files changed, 50 insertions(+), 61 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; diff --git a/intern/cycles/kernel/split/kernel_background_buffer_update.h b/intern/cycles/kernel/split/kernel_background_buffer_update.h index 181a1054a0d..0132ef9c2f2 100644 --- a/intern/cycles/kernel/split/kernel_background_buffer_update.h +++ b/intern/cycles/kernel/split/kernel_background_buffer_update.h @@ -161,15 +161,14 @@ ccl_device char kernel_background_buffer_update( ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER); } - if(IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) - { + if(IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) { #ifdef __BACKGROUND__ - /* sample background shader */ - float3 L_background = indirect_background(kg, state, ray, sd); - path_radiance_accum_background(L, (*throughput), L_background, state->bounce); + /* sample background shader */ + float3 L_background = indirect_background(kg, state, ray, sd); + path_radiance_accum_background(L, (*throughput), L_background, state->bounce); #endif - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER); - } + ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER); + } } if(IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER)) { diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h index 2cd98e466c1..421e2356f9b 100644 --- a/intern/cycles/kernel/split/kernel_data_init.h +++ b/intern/cycles/kernel/split/kernel_data_init.h @@ -291,13 +291,13 @@ ccl_device void kernel_data_init( /* Initialize queue data and queue index. */ if(thread_index < queuesize) { - /* Initialize active ray queue */ + /* Initialize active ray queue. */ Queue_data[QUEUE_ACTIVE_AND_REGENERATED_RAYS * queuesize + thread_index] = QUEUE_EMPTY_SLOT; - /* Initialize background and buffer update queue */ + /* Initialize background and buffer update queue. */ Queue_data[QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS * queuesize + thread_index] = QUEUE_EMPTY_SLOT; - /* Initialize shadow ray cast of AO queue */ + /* Initialize shadow ray cast of AO queue. */ Queue_data[QUEUE_SHADOW_RAY_CAST_AO_RAYS * queuesize + thread_index] = QUEUE_EMPTY_SLOT; - /* Initialize shadow ray cast of direct lighting queue */ + /* Initialize shadow ray cast of direct lighting queue. */ Queue_data[QUEUE_SHADOW_RAY_CAST_DL_RAYS * queuesize + thread_index] = QUEUE_EMPTY_SLOT; } @@ -316,7 +316,6 @@ ccl_device void kernel_data_init( int y = get_global_id(1); if(x < (sw * parallel_samples) && y < sh) { - int ray_index = x + y * (sw * parallel_samples); /* This is the first assignment to ray_state; @@ -400,12 +399,10 @@ ccl_device void kernel_data_init( #endif } else { /* These rays do not participate in path-iteration. */ - float4 L_rad = make_float4(0.0f, 0.0f, 0.0f, 0.0f); /* Accumulate result in output buffer. */ kernel_write_pass_float4(per_sample_output_buffers, my_sample, L_rad); path_rng_end(kg, rng_state, rng_coop[ray_index]); - ASSIGN_RAY_STATE(ray_state, ray_index, RAY_TO_REGENERATE); } } diff --git a/intern/cycles/kernel/split/kernel_next_iteration_setup.h b/intern/cycles/kernel/split/kernel_next_iteration_setup.h index 2dbdabc5fd3..e1a1577d7ae 100644 --- a/intern/cycles/kernel/split/kernel_next_iteration_setup.h +++ b/intern/cycles/kernel/split/kernel_next_iteration_setup.h @@ -127,14 +127,13 @@ ccl_device char kernel_next_iteration_setup( } if(IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - ccl_global float3 *throughput = &throughput_coop[ray_index]; ccl_global Ray *ray = &Ray_coop[ray_index]; ccl_global RNG* rng = &rng_coop[ray_index]; state = &PathState_coop[ray_index]; L = &PathRadiance_coop[ray_index]; - /* compute direct lighting and next bounce */ + /* Compute direct lighting and next bounce. */ if(!kernel_path_surface_bounce(kg, rng, sd, throughput, state, L, ray)) { ASSIGN_RAY_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER); enqueue_flag = 1; diff --git a/intern/cycles/kernel/split/kernel_scene_intersect.h b/intern/cycles/kernel/split/kernel_scene_intersect.h index 73aa005a496..7eb201ecf32 100644 --- a/intern/cycles/kernel/split/kernel_scene_intersect.h +++ b/intern/cycles/kernel/split/kernel_scene_intersect.h @@ -131,7 +131,8 @@ ccl_device void kernel_scene_intersect( if(!hit) { /* Change the state of rays that hit the background; * These rays undergo special processing in the - * background_bufferUpdate kernel*/ + * background_bufferUpdate kernel. + */ ASSIGN_RAY_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND); } } diff --git a/intern/cycles/kernel/split/kernel_shadow_blocked.h b/intern/cycles/kernel/split/kernel_shadow_blocked.h index 154ec53ffbb..28351c2b1ae 100644 --- a/intern/cycles/kernel/split/kernel_shadow_blocked.h +++ b/intern/cycles/kernel/split/kernel_shadow_blocked.h @@ -65,7 +65,7 @@ ccl_device void kernel_shadow_blocked( if(IS_FLAG(ray_state, ray_index, RAY_SHADOW_RAY_CAST_DL) || IS_FLAG(ray_state, ray_index, RAY_SHADOW_RAY_CAST_AO)) { - /* Load kernel global structure */ + /* Load kernel global structure. */ KernelGlobals *kg = (KernelGlobals *)globals; ShaderData *sd_shadow = (ShaderData *)shader_shadow;