diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-11-01 14:51:32 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-11-01 15:22:42 +0300 |
commit | dc9e0b819b814bbbbb38d4238c2fec21d9092437 (patch) | |
tree | 46b6150b93e17b4e40628643eb7babc09b2971e4 /intern/cycles/kernel/split | |
parent | 84e8b05e97522b9c12bd389c931601814307e5ff (diff) |
Cycles: Remove unused argument from the split kernel functions
Should be no functional changes, just simplifies operation with kernels.
Diffstat (limited to 'intern/cycles/kernel/split')
9 files changed, 13 insertions, 21 deletions
diff --git a/intern/cycles/kernel/split/kernel_background_buffer_update.h b/intern/cycles/kernel/split/kernel_background_buffer_update.h index 5d50415ba13..e02e55b5f18 100644 --- a/intern/cycles/kernel/split/kernel_background_buffer_update.h +++ b/intern/cycles/kernel/split/kernel_background_buffer_update.h @@ -43,7 +43,7 @@ * Queue_index (QUEUE_ACTIVE_AND_REGENERATED_RAYS) ------| |--- work_array * parallel_samples -------------------------------------| |--- PathState_coop * end_sample -------------------------------------------| |--- throughput_coop - * kg (globals + data) ----------------------------------| |--- rng_coop + * kg (globals) -----------------------------------------| |--- rng_coop * rng_state --------------------------------------------| |--- Ray * PathRadiance_coop ------------------------------------| | * sw ---------------------------------------------------| | @@ -71,7 +71,6 @@ */ ccl_device char kernel_background_buffer_update( KernelGlobals *kg, - ccl_constant KernelData *data, ShaderData *sd, ccl_global float *per_sample_output_buffers, ccl_global uint *rng_state, diff --git a/intern/cycles/kernel/split/kernel_direct_lighting.h b/intern/cycles/kernel/split/kernel_direct_lighting.h index f76e0cfaad4..34000517006 100644 --- a/intern/cycles/kernel/split/kernel_direct_lighting.h +++ b/intern/cycles/kernel/split/kernel_direct_lighting.h @@ -33,7 +33,7 @@ * sd -----------------------------------------------| |--- LightRay_coop * ray_state ----------------------------------------| |--- ray_state * Queue_data (QUEUE_ACTIVE_AND_REGENERATED_RAYS) ---| | - * kg (globals + data) ------------------------------| | + * kg (globals) -------------------------------------| | * queuesize ----------------------------------------| | * * note on sd_DL : sd_DL is neither input nor output to this kernel; sd_DL is filled and consumed in this kernel itself. @@ -50,7 +50,6 @@ */ ccl_device char kernel_direct_lighting( KernelGlobals *kg, - ccl_constant KernelData *data, ShaderData *sd, /* Required for direct lighting */ ShaderData *sd_DL, /* Required for direct lighting */ ccl_global uint *rng_coop, /* Required for 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 7681b95bb35..78dada82d89 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 @@ -40,7 +40,7 @@ * 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 - * kg (globals + data) ----------------------------------| |--- AOBSDF_coop + * kg (globals) -----------------------------------------| |--- AOBSDF_coop * parallel_samples -------------------------------------| |--- AOLightRay_coop * per_sample_output_buffers ----------------------------| | * sw ---------------------------------------------------| | @@ -72,7 +72,6 @@ */ ccl_device void kernel_holdout_emission_blurring_pathtermination_ao( KernelGlobals *kg, - ccl_constant KernelData *data, 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 */ diff --git a/intern/cycles/kernel/split/kernel_lamp_emission.h b/intern/cycles/kernel/split/kernel_lamp_emission.h index f5143f595a7..53dd6621127 100644 --- a/intern/cycles/kernel/split/kernel_lamp_emission.h +++ b/intern/cycles/kernel/split/kernel_lamp_emission.h @@ -26,7 +26,7 @@ * Throughput_coop ------------------------------------|--- kernel_lamp_emission --|--- PathRadiance_coop * Ray_coop -------------------------------------------| |--- Queue_data(QUEUE_ACTIVE_AND_REGENERATED_RAYS) * PathState_coop -------------------------------------| |--- Queue_index(QUEUE_ACTIVE_AND_REGENERATED_RAYS) - * kg (globals + data) --------------------------------| | + * kg (globals) ---------------------------------------| | * Intersection_coop ----------------------------------| | * ray_state ------------------------------------------| | * Queue_data (QUEUE_ACTIVE_AND_REGENERATED_RAYS) -----| | @@ -41,7 +41,6 @@ */ ccl_device void kernel_lamp_emission( KernelGlobals *kg, - ccl_constant KernelData *data, ShaderData *sd, /* Required for lamp emission */ ccl_global float3 *throughput_coop, /* Required for lamp emission */ PathRadiance *PathRadiance_coop, /* Required for lamp emission */ diff --git a/intern/cycles/kernel/split/kernel_next_iteration_setup.h b/intern/cycles/kernel/split/kernel_next_iteration_setup.h index fa8793203d1..1f1d7fb43f4 100644 --- a/intern/cycles/kernel/split/kernel_next_iteration_setup.h +++ b/intern/cycles/kernel/split/kernel_next_iteration_setup.h @@ -35,7 +35,7 @@ * Queue_data (QUEUE_ACTIVE_AND_REGENERATD_RAYS) --------| |--- Ray_coop * Queue_index (QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS) ---| |--- use_queues_flag * Ray_coop ---------------------------------------------| | - * kg (globals + data) ----------------------------------| | + * kg (globals) -----------------------------------------| | * LightRay_dl_coop -------------------------------------| * ISLamp_coop ------------------------------------------| * BSDFEval_coop ----------------------------------------| @@ -61,7 +61,6 @@ */ ccl_device char kernel_next_iteration_setup( KernelGlobals *kg, - ccl_constant KernelData *data, 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 */ diff --git a/intern/cycles/kernel/split/kernel_scene_intersect.h b/intern/cycles/kernel/split/kernel_scene_intersect.h index 1871f5d58cf..1a8b849fad0 100644 --- a/intern/cycles/kernel/split/kernel_scene_intersect.h +++ b/intern/cycles/kernel/split/kernel_scene_intersect.h @@ -32,7 +32,7 @@ * use_queues_flag --------------------------------| | * parallel_samples -------------------------------| | * QueueData(QUEUE_ACTIVE_AND_REGENERATED_RAYS) ---| | - * kg (data + globals) ----------------------------| | + * kg (globals) -----------------------------------| | * rng_coop ---------------------------------------| | * sw ---------------------------------------------| | * sh ---------------------------------------------| | @@ -64,7 +64,6 @@ ccl_device void kernel_scene_intersect( KernelGlobals *kg, - ccl_constant KernelData *data, ccl_global uint *rng_coop, ccl_global Ray *Ray_coop, /* Required for scene_intersect */ ccl_global PathState *PathState_coop, /* Required for scene_intersect */ diff --git a/intern/cycles/kernel/split/kernel_shader_eval.h b/intern/cycles/kernel/split/kernel_shader_eval.h index da6b2229543..dca7d67ee93 100644 --- a/intern/cycles/kernel/split/kernel_shader_eval.h +++ b/intern/cycles/kernel/split/kernel_shader_eval.h @@ -30,7 +30,7 @@ * Queue_data (QUEUE_ACTIVE_AND_REGENERATD_RAYS)-------| | * Queue_index(QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS)---| | * ray_state ------------------------------------------| | - * kg (globals + data) --------------------------------| | + * kg (globals) ---------------------------------------| | * queuesize ------------------------------------------| | * * Note on Queues : @@ -46,7 +46,6 @@ */ ccl_device void kernel_shader_eval( KernelGlobals *kg, - ccl_constant KernelData *data, 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 */ diff --git a/intern/cycles/kernel/split/kernel_shadow_blocked.h b/intern/cycles/kernel/split/kernel_shadow_blocked.h index 290424e2ecb..5f515a98783 100644 --- a/intern/cycles/kernel/split/kernel_shadow_blocked.h +++ b/intern/cycles/kernel/split/kernel_shadow_blocked.h @@ -31,7 +31,7 @@ QUEUE_SHADOW_RAY_CAST_DL_RAYS) -------| | * Queue_index(QUEUE_SHADOW_RAY_CAST_AO_RAYS& QUEUE_SHADOW_RAY_CAST_DL_RAYS) -------| | - * kg (globals + data) -----------------------------| | + * kg (globals) ------------------------------------| | * queuesize ---------------------------------------| | * * Note on sd_shadow : sd_shadow is neither input nor output to this kernel. sd_shadow is filled and consumed in this kernel itself. @@ -47,7 +47,6 @@ */ ccl_device void kernel_shadow_blocked( KernelGlobals *kg, - ccl_constant KernelData *data, 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 */ diff --git a/intern/cycles/kernel/split/kernel_sum_all_radiance.h b/intern/cycles/kernel/split/kernel_sum_all_radiance.h index a21e9b6a0b1..c860fc42a86 100644 --- a/intern/cycles/kernel/split/kernel_sum_all_radiance.h +++ b/intern/cycles/kernel/split/kernel_sum_all_radiance.h @@ -24,7 +24,7 @@ * by all different samples and stores them in the RenderTile's output buffer. */ ccl_device void kernel_sum_all_radiance( - ccl_constant KernelData *data, /* To get pass_stride to offet into buffer */ + KernelGlobals *kg, /* To get pass_stride to offet into buffer */ ccl_global float *buffer, /* Output buffer of RenderTile */ ccl_global float *per_sample_output_buffer, /* Radiance contributed by all samples */ int parallel_samples, int sw, int sh, int stride, @@ -37,14 +37,14 @@ ccl_device void kernel_sum_all_radiance( int y = get_global_id(1); if(x < sw && y < sh) { - buffer += ((buffer_offset_x + x) + (buffer_offset_y + y) * buffer_stride) * (data->film.pass_stride); - per_sample_output_buffer += ((x + y * stride) * parallel_samples) * (data->film.pass_stride); + buffer += ((buffer_offset_x + x) + (buffer_offset_y + y) * buffer_stride) * (kg->data->film.pass_stride); + per_sample_output_buffer += ((x + y * stride) * parallel_samples) * (kg->data->film.pass_stride); - int sample_stride = (data->film.pass_stride); + int sample_stride = (kg->data->film.pass_stride); int sample_iterator = 0; int pass_stride_iterator = 0; - int num_floats = data->film.pass_stride; + int num_floats = kg->data->film.pass_stride; for(sample_iterator = 0; sample_iterator < parallel_samples; sample_iterator++) { for(pass_stride_iterator = 0; pass_stride_iterator < num_floats; pass_stride_iterator++) { |