Cycles: Partially revert previous commit

Apparently removing kernel arguments broke NVidia OpenCL.

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

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++) {