Cycles: Code cleanup in split kernel, whitespaces

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;