Cycles: Remove unused argument from the split kernel functions

Should be no functional changes, just simplifies operation with kernels.
This commit is contained in:
Sergey Sharybin 2015-11-01 16:51:32 +05:00
parent 84e8b05e97
commit dc9e0b819b
19 changed files with 16 additions and 48 deletions

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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,

View File

@ -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 */

View File

@ -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 */

View File

@ -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 */

View File

@ -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 */

View File

@ -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 */

View File

@ -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 */

View File

@ -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 */

View File

@ -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++) {