Cycles: Remove unused argument from the split kernel functions
Should be no functional changes, just simplifies operation with kernels.
This commit is contained in:
parent
84e8b05e97
commit
dc9e0b819b
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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,
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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 */
|
||||
|
|
|
@ -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++) {
|
||||
|
|
Loading…
Reference in New Issue