blender
/
blender
129
397
Fork 571
Code Issues 6.2k Pull Requests 751 Projects 19 Wiki Activity

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:
Sergey Sharybin 2015-11-01 21:00:26 +05:00
parent dc9e0b819b
commit 9bce104c8c
11 changed files with 24 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
intern/cycles/device/device_opencl.cpp
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_holdout_emission_blurring_pathtermination_ao.cl
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
View File

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

1
intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
View File

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

4
intern/cycles/kernel/kernels/opencl/kernel_sum_all_radiance.cl
View File

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

10
intern/cycles/kernel/split/kernel_sum_all_radiance.h
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(
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++) {