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 | |
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')
19 files changed, 16 insertions, 48 deletions
diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp index d6b1ab38bd0..b4fc41ab247 100644 --- a/intern/cycles/device/device_opencl.cpp +++ b/intern/cycles/device/device_opencl.cpp @@ -2843,7 +2843,6 @@ public: kernel_set_args(ckPathTraceKernel_scene_intersect, 0, kgbuffer, - d_data, rng_coop, Ray_coop, PathState_coop, @@ -2863,7 +2862,6 @@ public: kernel_set_args(ckPathTraceKernel_lamp_emission, 0, kgbuffer, - d_data, sd, throughput_coop, PathRadiance_coop, @@ -2889,7 +2887,6 @@ public: kernel_set_args(ckPathTraceKernel_background_buffer_update, 0, kgbuffer, - d_data, sd, per_sample_output_buffers, d_rng_state, @@ -2926,7 +2923,6 @@ public: kernel_set_args(ckPathTraceKernel_shader_eval, 0, kgbuffer, - d_data, sd, rng_coop, Ray_coop, @@ -2940,7 +2936,6 @@ public: kernel_set_args(ckPathTraceKernel_holdout_emission_blurring_pathtermination_ao, 0, kgbuffer, - d_data, sd, per_sample_output_buffers, rng_coop, @@ -2970,7 +2965,6 @@ public: kernel_set_args(ckPathTraceKernel_direct_lighting, 0, kgbuffer, - d_data, sd, sd_DL_shadow, rng_coop, @@ -2986,7 +2980,6 @@ public: kernel_set_args(ckPathTraceKernel_shadow_blocked, 0, kgbuffer, - d_data, sd_DL_shadow, PathState_coop, LightRay_coop, @@ -3002,7 +2995,6 @@ public: kernel_set_args(ckPathTraceKernel_next_iteration_setup, 0, kgbuffer, - d_data, sd, rng_coop, throughput_coop, @@ -3023,7 +3015,7 @@ public: kernel_set_args(ckPathTraceKernel_sum_all_radiance, 0, - d_data, + kgbuffer, d_buffer, per_sample_output_buffers, num_parallel_samples, 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 c5eb7a24f76..c28338be5f2 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_background_buffer_update( ccl_global char *kg, - ccl_constant KernelData *data, ccl_global char *sd, ccl_global float *per_sample_output_buffers, ccl_global uint *rng_state, @@ -84,7 +83,6 @@ __kernel void kernel_ocl_path_trace_background_buffer_update( #endif enqueue_flag = kernel_background_buffer_update((KernelGlobals *)kg, - data, (ShaderData *)sd, per_sample_output_buffers, rng_state, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl index d23e87e8fd6..a646d42193d 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_direct_lighting( ccl_global char *kg, - ccl_constant KernelData *data, 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 */ @@ -62,7 +61,6 @@ __kernel void kernel_ocl_path_trace_direct_lighting( if(ray_index != QUEUE_EMPTY_SLOT) { #endif enqueue_flag = kernel_direct_lighting((KernelGlobals *)kg, - data, (ShaderData *)sd, (ShaderData *)sd_DL, rng_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 6fe25d7e48d..2f6d3e2d404 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 @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao( ccl_global char *kg, - ccl_constant KernelData *data, 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 */ @@ -76,7 +75,6 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao #endif kernel_holdout_emission_blurring_pathtermination_ao( (KernelGlobals *)kg, - data, (ShaderData *)sd, per_sample_output_buffers, rng_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl index feee4b7a01e..73de387ba86 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_lamp_emission( ccl_global char *kg, - ccl_constant KernelData *data, ccl_global char *sd, /* Required for lamp emission */ ccl_global float3 *throughput_coop, /* Required for lamp emission */ PathRadiance *PathRadiance_coop, /* Required for lamp emission */ @@ -69,7 +68,6 @@ __kernel void kernel_ocl_path_trace_lamp_emission( } kernel_lamp_emission((KernelGlobals *)kg, - data, (ShaderData *)sd, throughput_coop, PathRadiance_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 8896fe739d1..f6962aad9e8 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_next_iteration_setup( ccl_global char *kg, - ccl_constant KernelData *data, 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 */ @@ -84,7 +83,6 @@ __kernel void kernel_ocl_path_trace_next_iteration_setup( if(ray_index != QUEUE_EMPTY_SLOT) { #endif enqueue_flag = kernel_next_iteration_setup((KernelGlobals *)kg, - data, (ShaderData *)sd, rng_coop, throughput_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl index c9bbd65c706..c117008b99c 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_scene_intersect( ccl_global char *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 */ @@ -66,7 +65,6 @@ __kernel void kernel_ocl_path_trace_scene_intersect( } kernel_scene_intersect((KernelGlobals *)kg, - data, rng_coop, Ray_coop, PathState_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl index 2fa5496a84b..6f9faf509ac 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_shader_eval( ccl_global char *kg, - ccl_constant KernelData *data, 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 */ @@ -58,7 +57,6 @@ __kernel void kernel_ocl_path_trace_shader_eval( /* Continue on with shader evaluation. */ kernel_shader_eval((KernelGlobals *)kg, - data, (ShaderData *)sd, rng_coop, Ray_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl index b364c769840..b2ee6d2b776 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl @@ -18,7 +18,6 @@ __kernel void kernel_ocl_path_trace_shadow_blocked( ccl_global char *kg, - ccl_constant KernelData *data, 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 */ @@ -69,7 +68,6 @@ __kernel void kernel_ocl_path_trace_shadow_blocked( return; kernel_shadow_blocked((KernelGlobals *)kg, - data, (ShaderData *)sd_shadow, PathState_coop, LightRay_dl_coop, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl b/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl index 88a1ed830af..0ebfae911a8 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl @@ -17,7 +17,7 @@ #include "split/kernel_sum_all_radiance.h" __kernel void kernel_ocl_path_trace_sum_all_radiance( - ccl_constant KernelData *data, /* To get pass_stride to offet into buffer */ + ccl_global char *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, @@ -26,7 +26,7 @@ __kernel void kernel_ocl_path_trace_sum_all_radiance( int buffer_stride, int start_sample) { - kernel_sum_all_radiance(data, + kernel_sum_all_radiance((KernelGlobals *)kg, buffer, per_sample_output_buffer, parallel_samples, 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++) { |