diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-11-01 19:00:26 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2015-11-01 19:01:12 +0300 |
commit | 9bce104c8c72cdd4deac8e59ce571892302fb366 (patch) | |
tree | 91268823671c81e2df13af86fd2032c8b8725106 /intern | |
parent | dc9e0b819b814bbbbb38d4238c2fec21d9092437 (diff) |
Cycles: Partially revert previous commit
Apparently removing kernel arguments broke NVidia OpenCL.
Needs more investigation, for the time being revering changes which caused problem.
Diffstat (limited to 'intern')
11 files changed, 24 insertions, 8 deletions
diff --git a/intern/cycles/device/device_opencl.cpp b/intern/cycles/device/device_opencl.cpp index b4fc41ab247..d6b1ab38bd0 100644 --- a/intern/cycles/device/device_opencl.cpp +++ b/intern/cycles/device/device_opencl.cpp @@ -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, 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 c28338be5f2..a3eecd3128b 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl @@ -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, diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl index a646d42193d..d4a7cffb403 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl @@ -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 */ 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 2f6d3e2d404..e063614da1a 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,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 */ diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl index 73de387ba86..5215a0e0827 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl @@ -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 */ 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 f6962aad9e8..6d49b6294a8 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl @@ -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 */ diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl index c117008b99c..12eff6ccc1a 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl @@ -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 */ diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl index 6f9faf509ac..c37856c8f30 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl @@ -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 */ diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl index b2ee6d2b776..260a601946f 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl @@ -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 */ 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 0ebfae911a8..88a1ed830af 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_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, diff --git a/intern/cycles/kernel/split/kernel_sum_all_radiance.h b/intern/cycles/kernel/split/kernel_sum_all_radiance.h index c860fc42a86..a21e9b6a0b1 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( - 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++) { |