diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-10-29 19:44:36 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-10-29 19:52:56 +0300 |
commit | 4ca688a963408fe326ba8e2a696d7b0fdc4eb1e4 (patch) | |
tree | eeffd8dbbb5bf80dd98220cc3d501749c030c6e6 /intern | |
parent | fc5f717888f11caaa9cd246e2131a3892c81fbd1 (diff) |
Cycles: OpenCL split kernel cleanup, move casts from .h files to .cl files
Ideally we shouldn't use char* at all, but for now we have to, so at least
let's assume common .h files are free from pointer magic.
Diffstat (limited to 'intern')
18 files changed, 68 insertions, 103 deletions
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl index eff77b89a0a..c5eb7a24f76 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl @@ -17,9 +17,9 @@ #include "split/kernel_background_buffer_update.h" __kernel void kernel_ocl_path_trace_background_buffer_update( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, + ccl_global char *sd, ccl_global float *per_sample_output_buffers, ccl_global uint *rng_state, ccl_global uint *rng_coop, /* Required for buffer Update */ @@ -83,9 +83,9 @@ __kernel void kernel_ocl_path_trace_background_buffer_update( if(ray_index != QUEUE_EMPTY_SLOT) { #endif enqueue_flag = - kernel_background_buffer_update(globals, + kernel_background_buffer_update((KernelGlobals *)kg, data, - shader_data, + (ShaderData *)sd, per_sample_output_buffers, rng_state, rng_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl index c3277676029..8329aa98d35 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl @@ -17,9 +17,9 @@ #include "split/kernel_data_init.h" __kernel void kernel_ocl_path_trace_data_init( - ccl_global char *globals, - ccl_global char *shader_data_sd, /* Arguments related to ShaderData */ - ccl_global char *shader_data_sd_DL_shadow, /* Arguments related to ShaderData */ + ccl_global char *kg, + ccl_global char *sd, + ccl_global char *sd_DL_shadow, ccl_global float3 *P_sd, ccl_global float3 *P_sd_DL_shadow, @@ -141,9 +141,9 @@ __kernel void kernel_ocl_path_trace_data_init( #endif int parallel_samples) /* Number of samples to be processed in parallel */ { - kernel_data_init(globals, - shader_data_sd, - shader_data_sd_DL_shadow, + kernel_data_init((KernelGlobals *)kg, + (ShaderData *)sd, + (ShaderData *)sd_DL_shadow, P_sd, P_sd_DL_shadow, N_sd, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl index 6ec75013b3a..83e0afdbdeb 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl @@ -17,10 +17,10 @@ #include "split/kernel_direct_lighting.h" __kernel void kernel_ocl_path_trace_direct_lighting( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required for direct lighting */ - ccl_global char *shader_DL, /* Required for direct lighting */ + ccl_global char *sd, /* Required for direct lighting */ + ccl_global char *sd_DL, /* Required for direct lighting */ ccl_global uint *rng_coop, /* Required for direct lighting */ ccl_global PathState *PathState_coop, /* Required for direct lighting */ ccl_global int *ISLamp_coop, /* Required for direct lighting */ @@ -61,10 +61,10 @@ __kernel void kernel_ocl_path_trace_direct_lighting( #ifndef __COMPUTE_DEVICE_GPU__ if(ray_index != QUEUE_EMPTY_SLOT) { #endif - enqueue_flag = kernel_direct_lighting(globals, + enqueue_flag = kernel_direct_lighting((KernelGLobals *)kg, data, - shader_data, - shader_DL, + (ShaderData *)sd, + (ShaderData *)sd_DL, rng_coop, PathState_coop, ISLamp_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl index ae5f5cd1b3b..6fe25d7e48d 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl @@ -17,9 +17,9 @@ #include "split/kernel_holdout_emission_blurring_pathtermination_ao.h" __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required throughout the kernel except probabilistic path termination and AO */ + ccl_global char *sd, /* Required throughout the kernel except probabilistic path termination and AO */ ccl_global float *per_sample_output_buffers, ccl_global uint *rng_coop, /* Required for "kernel_write_data_passes" and AO */ ccl_global float3 *throughput_coop, /* Required for handling holdout material and AO */ @@ -75,9 +75,9 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao if(ray_index != QUEUE_EMPTY_SLOT) { #endif kernel_holdout_emission_blurring_pathtermination_ao( - globals, + (KernelGlobals *)kg, data, - shader_data, + (ShaderData *)sd, per_sample_output_buffers, rng_coop, throughput_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl index 1bc7808d834..fb327a32247 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl @@ -17,9 +17,9 @@ #include "split/kernel_lamp_emission.h" __kernel void kernel_ocl_path_trace_lamp_emission( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required for lamp emission */ + ccl_global char *sd, /* 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 */ @@ -68,9 +68,9 @@ __kernel void kernel_ocl_path_trace_lamp_emission( } } - kernel_lamp_emission(globals, + kernel_lamp_emission((KenrelGLobals *)kg, data, - shader_data, + (ShaderData *)sd, throughput_coop, PathRadiance_coop, Ray_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl index dcf4db40411..8896fe739d1 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl @@ -17,9 +17,9 @@ #include "split/kernel_next_iteration_setup.h" __kernel void kernel_ocl_path_trace_next_iteration_setup( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required for setting up ray for next iteration */ + ccl_global char *sd, /* Required for setting up ray for next iteration */ ccl_global uint *rng_coop, /* Required for setting up ray for next iteration */ ccl_global float3 *throughput_coop, /* Required for setting up ray for next iteration */ PathRadiance *PathRadiance_coop, /* Required for setting up ray for next iteration */ @@ -83,9 +83,9 @@ __kernel void kernel_ocl_path_trace_next_iteration_setup( #ifndef __COMPUTE_DEVICE_GPU__ if(ray_index != QUEUE_EMPTY_SLOT) { #endif - enqueue_flag = kernel_next_iteration_setup(globals, + enqueue_flag = kernel_next_iteration_setup((KernelGlobals *)kg, data, - shader_data, + (ShaderData *)sd, rng_coop, throughput_coop, PathRadiance_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl index e5fad7bce50..3750a0d0062 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl @@ -17,7 +17,7 @@ #include "split/kernel_scene_intersect.h" __kernel void kernel_ocl_path_trace_scene_intersect( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, ccl_global uint *rng_coop, ccl_global Ray *Ray_coop, /* Required for scene_intersect */ @@ -65,7 +65,7 @@ __kernel void kernel_ocl_path_trace_scene_intersect( } } - kernel_scene_intersect(globals, + kernel_scene_intersect((KernelGlobals *)kg, data, rng_coop, Ray_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl index b9f616e6bdf..2fa5496a84b 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl @@ -17,9 +17,9 @@ #include "split/kernel_shader_eval.h" __kernel void kernel_ocl_path_trace_shader_eval( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Output ShaderData structure to be filled */ + ccl_global char *sd, /* Output ShaderData structure to be filled */ ccl_global uint *rng_coop, /* Required for rbsdf calculation */ ccl_global Ray *Ray_coop, /* Required for setting up shader from ray */ ccl_global PathState *PathState_coop, /* Required for all functions in this kernel */ @@ -57,9 +57,9 @@ __kernel void kernel_ocl_path_trace_shader_eval( Queue_index); /* Continue on with shader evaluation. */ - kernel_shader_eval(globals, + kernel_shader_eval((KernelGlobals *)kg, data, - shader_data, + (ShaderData *)sd, rng_coop, Ray_coop, PathState_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl index 03886c0a030..b364c769840 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl @@ -17,9 +17,9 @@ #include "split/kernel_shadow_blocked.h" __kernel void kernel_ocl_path_trace_shadow_blocked( - ccl_global char *globals, + ccl_global char *kg, ccl_constant KernelData *data, - ccl_global char *shader_shadow, /* Required for shadow blocked */ + ccl_global char *sd_shadow, /* Required for shadow blocked */ ccl_global PathState *PathState_coop, /* Required for shadow blocked */ ccl_global Ray *LightRay_dl_coop, /* Required for direct lighting's shadow blocked */ ccl_global Ray *LightRay_ao_coop, /* Required for AO's shadow blocked */ @@ -68,9 +68,9 @@ __kernel void kernel_ocl_path_trace_shadow_blocked( if(ray_index == QUEUE_EMPTY_SLOT) return; - kernel_shadow_blocked(globals, + kernel_shadow_blocked((KernelGlobals *)kg, data, - shader_shadow, + (ShaderData *)sd_shadow, PathState_coop, LightRay_dl_coop, LightRay_ao_coop, diff --git a/intern/cycles/kernel/split/kernel_background_buffer_update.h b/intern/cycles/kernel/split/kernel_background_buffer_update.h index 0132ef9c2f2..5d50415ba13 100644 --- a/intern/cycles/kernel/split/kernel_background_buffer_update.h +++ b/intern/cycles/kernel/split/kernel_background_buffer_update.h @@ -57,7 +57,7 @@ * work_pool_wgs ----------------------------------------| | * num_samples ------------------------------------------| | * - * note on shader_data : shader_data argument is neither an input nor an output for this kernel. It is just filled and consumed here itself. + * note on sd : sd argument is neither an input nor an output for this kernel. It is just filled and consumed here itself. * Note on Queues : * This kernel fetches rays from QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue. * @@ -70,9 +70,9 @@ * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be empty */ ccl_device char kernel_background_buffer_update( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, + ShaderData *sd, ccl_global float *per_sample_output_buffers, ccl_global uint *rng_state, ccl_global uint *rng_coop, /* Required for buffer Update */ @@ -100,11 +100,6 @@ ccl_device char kernel_background_buffer_update( int ray_index) { char enqueue_flag = 0; - - /* Load kernel globals structure and ShaderData strucuture */ - KernelGlobals *kg = (KernelGlobals *)globals; - ShaderData *sd = (ShaderData *)shader_data; - #ifdef __KERNEL_DEBUG__ DebugData *debug_data = &debugdata_coop[ray_index]; #endif diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h index 4dab79a5c67..a5db6a83283 100644 --- a/intern/cycles/kernel/split/kernel_data_init.h +++ b/intern/cycles/kernel/split/kernel_data_init.h @@ -51,9 +51,9 @@ * The number of elements in the queues is initialized to 0; */ ccl_device void kernel_data_init( - ccl_global char *globals, - ccl_global char *shader_data_sd, /* Arguments related to ShaderData */ - ccl_global char *shader_data_sd_DL_shadow, /* Arguments related to ShaderData */ + KernelGlobals *kg, + ShaderData *sd, + ShaderData *sd_DL_shadow, ccl_global float3 *P_sd, ccl_global float3 *P_sd_DL_shadow, @@ -175,19 +175,11 @@ ccl_device void kernel_data_init( #endif int parallel_samples) /* Number of samples to be processed in parallel */ { - - /* Load kernel globals structure */ - KernelGlobals *kg = (KernelGlobals *)globals; - kg->data = data; #define KERNEL_TEX(type, ttype, name) \ kg->name = name; #include "../kernel_textures.h" - /* Load ShaderData structure */ - ShaderData *sd = (ShaderData *)shader_data_sd; - ShaderData *sd_DL_shadow = (ShaderData *)shader_data_sd_DL_shadow; - sd->P = P_sd; sd_DL_shadow->P = P_sd_DL_shadow; diff --git a/intern/cycles/kernel/split/kernel_direct_lighting.h b/intern/cycles/kernel/split/kernel_direct_lighting.h index 50c83d06140..f76e0cfaad4 100644 --- a/intern/cycles/kernel/split/kernel_direct_lighting.h +++ b/intern/cycles/kernel/split/kernel_direct_lighting.h @@ -30,13 +30,13 @@ * * rng_coop -----------------------------------------|--- kernel_direct_lighting --|--- BSDFEval_coop * PathState_coop -----------------------------------| |--- ISLamp_coop - * shader_data --------------------------------------| |--- LightRay_coop + * sd -----------------------------------------------| |--- LightRay_coop * ray_state ----------------------------------------| |--- ray_state * Queue_data (QUEUE_ACTIVE_AND_REGENERATED_RAYS) ---| | * kg (globals + data) ------------------------------| | * queuesize ----------------------------------------| | * - * note on shader_DL : shader_DL is neither input nor output to this kernel; shader_DL is filled and consumed in this kernel itself. + * note on sd_DL : sd_DL is neither input nor output to this kernel; sd_DL is filled and consumed in this kernel itself. * Note on Queues : * This kernel only reads from the QUEUE_ACTIVE_AND_REGENERATED_RAYS queue and processes * only the rays of state RAY_ACTIVE; If a ray needs to execute the corresponding shadow_blocked @@ -49,10 +49,10 @@ * kernel call. Before this kernel call the QUEUE_SHADOW_RAY_CAST_DL_RAYS will be empty. */ ccl_device char kernel_direct_lighting( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required for direct lighting */ - ccl_global char *shader_DL, /* Required for direct lighting */ + ShaderData *sd, /* Required for direct lighting */ + ShaderData *sd_DL, /* Required for direct lighting */ ccl_global uint *rng_coop, /* Required for direct lighting */ ccl_global PathState *PathState_coop, /* Required for direct lighting */ ccl_global int *ISLamp_coop, /* Required for direct lighting */ @@ -63,11 +63,6 @@ ccl_device char kernel_direct_lighting( { char enqueue_flag = 0; if(IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - /* Load kernel globals structure and ShaderData structure. */ - KernelGlobals *kg = (KernelGlobals *)globals; - ShaderData *sd = (ShaderData *)shader_data; - ShaderData *sd_DL = (ShaderData *)shader_DL; - ccl_global PathState *state = &PathState_coop[ray_index]; /* direct lighting */ diff --git a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h b/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h index a75523a3e53..7681b95bb35 100644 --- a/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h +++ b/intern/cycles/kernel/split/kernel_holdout_emission_blurring_pathtermination_ao.h @@ -36,7 +36,7 @@ * Intersection_coop ------------------------------------| |--- L_transparent_coop * PathState_coop ---------------------------------------| |--- per_sample_output_buffers * L_transparent_coop -----------------------------------| |--- PathRadiance_coop - * shader_data ------------------------------------------| |--- ShaderData + * sd ---------------------------------------------------| |--- ShaderData * ray_state --------------------------------------------| |--- ray_state * Queue_data (QUEUE_ACTIVE_AND_REGENERATED_RAYS) -------| |--- Queue_data (QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS) * Queue_index (QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS) ---| |--- AOAlpha_coop @@ -71,9 +71,9 @@ * QUEUE_SHADOW_RAY_CAST_AO_RAYS will be filled with rays marked with flag RAY_SHADOW_RAY_CAST_AO */ ccl_device void kernel_holdout_emission_blurring_pathtermination_ao( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required throughout the kernel except probabilistic path termination and AO */ + ShaderData *sd, /* Required throughout the kernel except probabilistic path termination and AO */ ccl_global float *per_sample_output_buffers, ccl_global uint *rng_coop, /* Required for "kernel_write_data_passes" and AO */ ccl_global float3 *throughput_coop, /* Required for handling holdout material and AO */ @@ -95,10 +95,6 @@ ccl_device void kernel_holdout_emission_blurring_pathtermination_ao( char *enqueue_flag, char *enqueue_flag_AO_SHADOW_RAY_CAST) { - /* Load kernel globals structure and ShaderData structure */ - KernelGlobals *kg = (KernelGlobals *)globals; - ShaderData *sd = (ShaderData *)shader_data; - #ifdef __WORK_STEALING__ unsigned int my_work; unsigned int pixel_x; diff --git a/intern/cycles/kernel/split/kernel_lamp_emission.h b/intern/cycles/kernel/split/kernel_lamp_emission.h index a8e4b0a06c8..f5143f595a7 100644 --- a/intern/cycles/kernel/split/kernel_lamp_emission.h +++ b/intern/cycles/kernel/split/kernel_lamp_emission.h @@ -37,12 +37,12 @@ * sh -------------------------------------------------| | * parallel_samples -----------------------------------| | * - * note : shader_data is neither input nor output. Its just filled and consumed in the same, kernel_lamp_emission, kernel. + * note : sd is neither input nor output. Its just filled and consumed in the same, kernel_lamp_emission, kernel. */ ccl_device void kernel_lamp_emission( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required for lamp emission */ + ShaderData *sd, /* 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 */ @@ -59,8 +59,6 @@ ccl_device void kernel_lamp_emission( 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; PathRadiance *L = &PathRadiance_coop[ray_index]; float3 throughput = throughput_coop[ray_index]; diff --git a/intern/cycles/kernel/split/kernel_next_iteration_setup.h b/intern/cycles/kernel/split/kernel_next_iteration_setup.h index e1a1577d7ae..fa8793203d1 100644 --- a/intern/cycles/kernel/split/kernel_next_iteration_setup.h +++ b/intern/cycles/kernel/split/kernel_next_iteration_setup.h @@ -30,7 +30,7 @@ * throughput_coop --------------------------------------| |--- Queue_data (QUEUE_HITBF_BUFF_UPDATE_TOREGEN_RAYS) * PathRadiance_coop ------------------------------------| |--- throughput_coop * PathState_coop ---------------------------------------| |--- PathRadiance_coop - * shader_data ------------------------------------------| |--- PathState_coop + * sd ---------------------------------------------------| |--- PathState_coop * ray_state --------------------------------------------| |--- ray_state * Queue_data (QUEUE_ACTIVE_AND_REGENERATD_RAYS) --------| |--- Ray_coop * Queue_index (QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS) ---| |--- use_queues_flag @@ -60,9 +60,9 @@ * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with RAY_TO_REGENERATE and more RAY_UPDATE_BUFFER rays */ ccl_device char kernel_next_iteration_setup( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Required for setting up ray for next iteration */ + ShaderData *sd, /* Required for setting up ray for next iteration */ ccl_global uint *rng_coop, /* Required for setting up ray for next iteration */ ccl_global float3 *throughput_coop, /* Required for setting up ray for next iteration */ PathRadiance *PathRadiance_coop, /* Required for setting up ray for next iteration */ @@ -81,11 +81,9 @@ ccl_device char kernel_next_iteration_setup( { char enqueue_flag = 0; - /* Load kernel globals structure and ShaderData structure. */ - KernelGlobals *kg = (KernelGlobals *)globals; - ShaderData *sd = (ShaderData *)shader_data; - PathRadiance *L = 0x0; - ccl_global PathState *state = 0x0; + /* Load ShaderData structure. */ + PathRadiance *L = NULL; + ccl_global PathState *state = NULL; /* Path radiance update for AO/Direct_lighting's shadow blocked. */ if(IS_FLAG(ray_state, ray_index, RAY_SHADOW_RAY_CAST_DL) || diff --git a/intern/cycles/kernel/split/kernel_scene_intersect.h b/intern/cycles/kernel/split/kernel_scene_intersect.h index 7eb201ecf32..1871f5d58cf 100644 --- a/intern/cycles/kernel/split/kernel_scene_intersect.h +++ b/intern/cycles/kernel/split/kernel_scene_intersect.h @@ -63,7 +63,7 @@ */ ccl_device void kernel_scene_intersect( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, ccl_global uint *rng_coop, ccl_global Ray *Ray_coop, /* Required for scene_intersect */ @@ -86,9 +86,6 @@ ccl_device void kernel_scene_intersect( if(!IS_STATE(ray_state, ray_index, RAY_ACTIVE)) return; - /* Load kernel globals structure */ - KernelGlobals *kg = (KernelGlobals *)globals; - #ifdef __KERNEL_DEBUG__ DebugData *debug_data = &debugdata_coop[ray_index]; #endif diff --git a/intern/cycles/kernel/split/kernel_shader_eval.h b/intern/cycles/kernel/split/kernel_shader_eval.h index e6fdc592586..da6b2229543 100644 --- a/intern/cycles/kernel/split/kernel_shader_eval.h +++ b/intern/cycles/kernel/split/kernel_shader_eval.h @@ -23,7 +23,7 @@ * the rays of state RAY_TO_REGENERATE and enqueues them in QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS queue. * * The input and output of the kernel is as follows, - * rng_coop -------------------------------------------|--- kernel_shader_eval --|--- shader_data + * rng_coop -------------------------------------------|--- kernel_shader_eval --|--- sd * Ray_coop -------------------------------------------| |--- Queue_data (QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS) * PathState_coop -------------------------------------| |--- Queue_index (QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS) * Intersection_coop ----------------------------------| | @@ -45,9 +45,9 @@ * QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS will be filled with RAY_TO_REGENERATE rays */ ccl_device void kernel_shader_eval( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, - ccl_global char *shader_data, /* Output ShaderData structure to be filled */ + ShaderData *sd, /* Output ShaderData structure to be filled */ ccl_global uint *rng_coop, /* Required for rbsdf calculation */ ccl_global Ray *Ray_coop, /* Required for setting up shader from ray */ ccl_global PathState *PathState_coop, /* Required for all functions in this kernel */ @@ -56,8 +56,6 @@ ccl_device void kernel_shader_eval( int ray_index) { if(IS_STATE(ray_state, ray_index, RAY_ACTIVE)) { - KernelGlobals *kg = (KernelGlobals *)globals; - ShaderData *sd = (ShaderData *)shader_data; Intersection *isect = &Intersection_coop[ray_index]; ccl_global uint *rng = &rng_coop[ray_index]; ccl_global PathState *state = &PathState_coop[ray_index]; diff --git a/intern/cycles/kernel/split/kernel_shadow_blocked.h b/intern/cycles/kernel/split/kernel_shadow_blocked.h index 28351c2b1ae..a3b9f23116c 100644 --- a/intern/cycles/kernel/split/kernel_shadow_blocked.h +++ b/intern/cycles/kernel/split/kernel_shadow_blocked.h @@ -34,7 +34,7 @@ * kg (globals + data) -----------------------------| | * queuesize ---------------------------------------| | * - * Note on shader_shadow : shader_shadow is neither input nor output to this kernel. shader_shadow is filled and consumed in this kernel itself. + * Note on sd_shadow : sd_shadow is neither input nor output to this kernel. sd_shadow is filled and consumed in this kernel itself. * Note on queues : * The kernel fetches from QUEUE_SHADOW_RAY_CAST_AO_RAYS and QUEUE_SHADOW_RAY_CAST_DL_RAYS queues. We will empty * these queues this kernel. @@ -46,9 +46,9 @@ * QUEUE_SHADOW_RAY_CAST_AO_RAYS and QUEUE_SHADOW_RAY_CAST_DL_RAYS will be empty at kernel exit. */ ccl_device void kernel_shadow_blocked( - ccl_global char *globals, + KernelGlobals *kg, ccl_constant KernelData *data, - ccl_global char *shader_shadow, /* Required for shadow blocked */ + ShaderData *sd_shadow, /* Required for shadow blocked */ ccl_global PathState *PathState_coop, /* Required for shadow blocked */ ccl_global Ray *LightRay_dl_coop, /* Required for direct lighting's shadow blocked */ ccl_global Ray *LightRay_ao_coop, /* Required for AO's shadow blocked */ @@ -65,10 +65,6 @@ 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. */ - KernelGlobals *kg = (KernelGlobals *)globals; - ShaderData *sd_shadow = (ShaderData *)shader_shadow; - ccl_global PathState *state = &PathState_coop[ray_index]; ccl_global Ray *light_ray_dl_global = &LightRay_dl_coop[ray_index]; ccl_global Ray *light_ray_ao_global = &LightRay_ao_coop[ray_index]; |