diff options
Diffstat (limited to 'intern/cycles/kernel/split/kernel_lamp_emission.h')
| -rw-r--r-- | intern/cycles/kernel/split/kernel_lamp_emission.h | 67 |
1 files changed, 17 insertions, 50 deletions
diff --git a/intern/cycles/kernel/split/kernel_lamp_emission.h b/intern/cycles/kernel/split/kernel_lamp_emission.h index b804bfc8630..e5fdb637a50 100644 --- a/intern/cycles/kernel/split/kernel_lamp_emission.h +++ b/intern/cycles/kernel/split/kernel_lamp_emission.h @@ -16,8 +16,7 @@ #include "kernel_split_common.h" -/* - * Note on kernel_lamp_emission +/* Note on kernel_lamp_emission * This is the 3rd kernel in the ray-tracing logic. This is the second of the * path-iteration kernels. This kernel takes care of the indirect lamp emission logic. * This kernel operates on QUEUE_ACTIVE_AND_REGENERATED_RAYS. It processes rays of state RAY_ACTIVE @@ -40,55 +39,23 @@ * * note : shader_data is neither input nor output. Its just filled and consumed in the same, kernel_lamp_emission, kernel. */ -__kernel void kernel_lamp_emission( - ccl_global char *globals, - ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required for lamp emission */ - ccl_global float3 *throughput_coop, /* Required for lamp emission */ - PathRadiance *PathRadiance_coop, /* Required for lamp emission */ - ccl_global Ray *Ray_coop, /* Required for lamp emission */ - ccl_global PathState *PathState_coop, /* Required for lamp emission */ - Intersection *Intersection_coop, /* Required for lamp emission */ - ccl_global char *ray_state, /* Denotes the state of each ray */ - int sw, int sh, - ccl_global int *Queue_data, /* Memory for queues */ - ccl_global int *Queue_index, /* Tracks the number of elements in queues */ - int queuesize, /* Size (capacity) of queues */ - ccl_global char *use_queues_flag, /* used to decide if this kernel should use queues to fetch ray index */ - int parallel_samples /* Number of samples to be processed in parallel */ - ) +ccl_device void kernel_lamp_emission( + ccl_global char *globals, + ccl_constant KernelData *data, + ccl_global char *shader_data, /* Required for lamp emission */ + ccl_global float3 *throughput_coop, /* Required for lamp emission */ + PathRadiance *PathRadiance_coop, /* Required for lamp emission */ + ccl_global Ray *Ray_coop, /* Required for lamp emission */ + ccl_global PathState *PathState_coop, /* Required for lamp emission */ + Intersection *Intersection_coop, /* Required for lamp emission */ + ccl_global char *ray_state, /* Denotes the state of each ray */ + int sw, int sh, + ccl_global char *use_queues_flag, /* Used to decide if this kernel should use + * queues to fetch ray index + */ + int parallel_samples, /* Number of samples to be processed in parallel */ + int ray_index) { - int x = get_global_id(0); - int y = get_global_id(1); - - /* We will empty this queue in this kernel */ - if(get_global_id(0) == 0 && get_global_id(1) == 0) { - Queue_index[QUEUE_ACTIVE_AND_REGENERATED_RAYS] = 0; - } - - /* Fetch use_queues_flag */ - ccl_local char local_use_queues_flag; - if(get_local_id(0) == 0 && get_local_id(1) == 0) { - local_use_queues_flag = use_queues_flag[0]; - } - barrier(CLK_LOCAL_MEM_FENCE); - - int ray_index; - if(local_use_queues_flag) { - int thread_index = get_global_id(1) * get_global_size(0) + get_global_id(0); - ray_index = get_ray_index(thread_index, QUEUE_ACTIVE_AND_REGENERATED_RAYS, Queue_data, queuesize, 1); - - if(ray_index == QUEUE_EMPTY_SLOT) { - return; - } - } else { - if(x < (sw * parallel_samples) && y < sh){ - ray_index = x + y * (sw * parallel_samples); - } else { - return; - } - } - if(IS_STATE(ray_state, ray_index, RAY_ACTIVE) || IS_STATE(ray_state, ray_index, RAY_HIT_BACKGROUND)) { KernelGlobals *kg = (KernelGlobals *)globals; ShaderData *sd = (ShaderData *)shader_data; |