forked from bartvdbraak/blender
Cycles: Partially revert previous commit
Apparently removing kernel arguments broke NVidia OpenCL. Needs more investigation, for the time being revering changes which caused problem.
This commit is contained in:
parent
dc9e0b819b
commit
9bce104c8c
@ -2843,6 +2843,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_scene_intersect,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
rng_coop,
|
||||
Ray_coop,
|
||||
PathState_coop,
|
||||
@ -2862,6 +2863,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_lamp_emission,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
sd,
|
||||
throughput_coop,
|
||||
PathRadiance_coop,
|
||||
@ -2887,6 +2889,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_background_buffer_update,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
sd,
|
||||
per_sample_output_buffers,
|
||||
d_rng_state,
|
||||
@ -2923,6 +2926,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_shader_eval,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
sd,
|
||||
rng_coop,
|
||||
Ray_coop,
|
||||
@ -2936,6 +2940,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_holdout_emission_blurring_pathtermination_ao,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
sd,
|
||||
per_sample_output_buffers,
|
||||
rng_coop,
|
||||
@ -2965,6 +2970,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_direct_lighting,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
sd,
|
||||
sd_DL_shadow,
|
||||
rng_coop,
|
||||
@ -2980,6 +2986,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_shadow_blocked,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
sd_DL_shadow,
|
||||
PathState_coop,
|
||||
LightRay_coop,
|
||||
@ -2995,6 +3002,7 @@ public:
|
||||
kernel_set_args(ckPathTraceKernel_next_iteration_setup,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
sd,
|
||||
rng_coop,
|
||||
throughput_coop,
|
||||
@ -3015,7 +3023,7 @@ public:
|
||||
|
||||
kernel_set_args(ckPathTraceKernel_sum_all_radiance,
|
||||
0,
|
||||
kgbuffer,
|
||||
d_data,
|
||||
d_buffer,
|
||||
per_sample_output_buffers,
|
||||
num_parallel_samples,
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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,
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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 */
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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 */
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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 */
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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 */
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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 */
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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 */
|
||||
|
@ -18,6 +18,7 @@
|
||||
|
||||
__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 */
|
||||
|
@ -17,7 +17,7 @@
|
||||
#include "split/kernel_sum_all_radiance.h"
|
||||
|
||||
__kernel void kernel_ocl_path_trace_sum_all_radiance(
|
||||
ccl_global char *kg, /* To get pass_stride to offet into buffer */
|
||||
ccl_constant KernelData *data, /* 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((KernelGlobals *)kg,
|
||||
kernel_sum_all_radiance(data,
|
||||
buffer,
|
||||
per_sample_output_buffer,
|
||||
parallel_samples,
|
||||
|
@ -24,7 +24,7 @@
|
||||
* by all different samples and stores them in the RenderTile's output buffer.
|
||||
*/
|
||||
ccl_device void kernel_sum_all_radiance(
|
||||
KernelGlobals *kg, /* To get pass_stride to offet into buffer */
|
||||
ccl_constant KernelData *data, /* 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) * (kg->data->film.pass_stride);
|
||||
per_sample_output_buffer += ((x + y * stride) * parallel_samples) * (kg->data->film.pass_stride);
|
||||
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);
|
||||
|
||||
int sample_stride = (kg->data->film.pass_stride);
|
||||
int sample_stride = (data->film.pass_stride);
|
||||
|
||||
int sample_iterator = 0;
|
||||
int pass_stride_iterator = 0;
|
||||
int num_floats = kg->data->film.pass_stride;
|
||||
int num_floats = 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
Block a user