Cycles: Code cleanup in split kernel, whitespaces
authorSergey Sharybin <sergey.vfx@gmail.com>
Fri, 3 Jul 2015 08:56:40 +0000 (10:56 +0200)
committerSergey Sharybin <sergey.vfx@gmail.com>
Fri, 3 Jul 2015 09:03:56 +0000 (11:03 +0200)
intern/cycles/kernel/kernel_queues.h
intern/cycles/kernel/split/kernel_background_buffer_update.h
intern/cycles/kernel/split/kernel_data_init.h
intern/cycles/kernel/split/kernel_next_iteration_setup.h
intern/cycles/kernel/split/kernel_scene_intersect.h
intern/cycles/kernel/split/kernel_shadow_blocked.h

index 9e65e2b0768036a5a43b2660ae9b5e66954478c6..cf5614b8a866ab23aed4218eb33f2b399ecb66b4 100644 (file)
 /*
  * 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;
index 181a1054a0d896877ddb73f4901a70f824d12af2..0132ef9c2f2b81e23ab0c242877eee32d1fcfc98 100644 (file)
@@ -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)) {
index 2cd98e466c1f45f57efbc29035938904b55441f2..421e2356f9b0b80e92e0ee633308315f920c28db 100644 (file)
@@ -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);
                }
        }
index 2dbdabc5fd3a2afe7864dc72aa268fba6017f250..e1a1577d7ae7940d72d02e7a1324229a92970195 100644 (file)
@@ -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;
index 73aa005a496c905c232168eb4ccdaf704046c4b9..7eb201ecf3209f80feb3a26ee2e552a7b092f7be 100644 (file)
@@ -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);
        }
 }
index 154ec53ffbb7ec2017a845d2365417c5a01e0269..28351c2b1ae1f4be2c870d3ffebb7ac287180852 100644 (file)
@@ -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;