Cycles: Partially revert previous commit

Apparently removing kernel arguments broke NVidia OpenCL.

Needs more investigation, for the time being revering changes which caused problem.

9 files changed, 10 insertions, 2 deletions

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,