[Bf-blender-cvs] [b9f89b1] master: Cycles: Code cleanup in split kernel, whitespaces
Sergey Sharybin
noreply at git.blender.org
Fri Jul 3 11:04:25 CEST 2015
Commit: b9f89b164785f28b023af071d696c298deba4f63
Author: Sergey Sharybin
Date: Fri Jul 3 10:56:40 2015 +0200
Branches: master
https://developer.blender.org/rBb9f89b164785f28b023af071d696c298deba4f63
Cycles: Code cleanup in split kernel, whitespaces
===================================================================
M intern/cycles/kernel/kernel_queues.h
M intern/cycles/kernel/split/kernel_background_buffer_update.h
M intern/cycles/kernel/split/kernel_data_init.h
M intern/cycles/kernel/split/kernel_next_iteration_setup.h
M intern/cycles/kernel/split/kernel_scene_intersect.h
M intern/cycles/kernel/split/kernel_shadow_blocked.h
===================================================================
diff --git a/intern/cycles/kernel/kernel_queues.h b/intern/cycles/kernel/kernel_queues.h
index 9e65e2b..cf5614b 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 181a105..0132ef9 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 2cd98e4..421e235 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 2dbdabc..e1a1577 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
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list