Cycles: OpenCL split kernel cleanup, move casts from .h files to .cl files

Ideally we shouldn't use char* at all, but for now we have to, so at least
let's assume common .h files are free from pointer magic.

9 files changed, 38 insertions, 38 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 eff77b89a0a..c5eb7a24f76 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_background_buffer_update.cl
@@ -17,9 +17,9 @@
 #include "split/kernel_background_buffer_update.h"
 
 __kernel void kernel_ocl_path_trace_background_buffer_update(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
-        ccl_global char *shader_data,
+        ccl_global char *sd,
         ccl_global float *per_sample_output_buffers,
         ccl_global uint *rng_state,
         ccl_global uint *rng_coop,             /* Required for buffer Update */
@@ -83,9 +83,9 @@ __kernel void kernel_ocl_path_trace_background_buffer_update(
 	if(ray_index != QUEUE_EMPTY_SLOT) {
 #endif
 		enqueue_flag =
-			kernel_background_buffer_update(globals,
+			kernel_background_buffer_update((KernelGlobals *)kg,
 			                                data,
-			                                shader_data,
+			                                (ShaderData *)sd,
 			                                per_sample_output_buffers,
 			                                rng_state,
 			                                rng_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
index c3277676029..8329aa98d35 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_data_init.cl
@@ -17,9 +17,9 @@
 #include "split/kernel_data_init.h"
 
 __kernel void kernel_ocl_path_trace_data_init(
-        ccl_global char *globals,
-        ccl_global char *shader_data_sd,                  /* Arguments related to ShaderData */
-        ccl_global char *shader_data_sd_DL_shadow,        /* Arguments related to ShaderData */
+        ccl_global char *kg,
+        ccl_global char *sd,
+        ccl_global char *sd_DL_shadow,
 
         ccl_global float3 *P_sd,
         ccl_global float3 *P_sd_DL_shadow,
@@ -141,9 +141,9 @@ __kernel void kernel_ocl_path_trace_data_init(
 #endif
         int parallel_samples)                        /* Number of samples to be processed in parallel */
 {
-	kernel_data_init(globals,
-	                 shader_data_sd,
-	                 shader_data_sd_DL_shadow,
+	kernel_data_init((KernelGlobals *)kg,
+	                 (ShaderData *)sd,
+	                 (ShaderData *)sd_DL_shadow,
 	                 P_sd,
 	                 P_sd_DL_shadow,
 	                 N_sd,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
index 6ec75013b3a..83e0afdbdeb 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_direct_lighting.cl
@@ -17,10 +17,10 @@
 #include "split/kernel_direct_lighting.h"
 
 __kernel void kernel_ocl_path_trace_direct_lighting(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
-        ccl_global char *shader_data,           /* Required for direct lighting */
-        ccl_global char *shader_DL,             /* Required for direct lighting */
+        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 */
         ccl_global PathState *PathState_coop,   /* Required for direct lighting */
         ccl_global int *ISLamp_coop,            /* Required for direct lighting */
@@ -61,10 +61,10 @@ __kernel void kernel_ocl_path_trace_direct_lighting(
 #ifndef __COMPUTE_DEVICE_GPU__
 	if(ray_index != QUEUE_EMPTY_SLOT) {
 #endif
-		enqueue_flag = kernel_direct_lighting(globals,
+		enqueue_flag = kernel_direct_lighting((KernelGLobals *)kg,
 		                                      data,
-		                                      shader_data,
-		                                      shader_DL,
+		                                      (ShaderData *)sd,
+		                                      (ShaderData *)sd_DL,
 		                                      rng_coop,
 		                                      PathState_coop,
 		                                      ISLamp_coop,
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 ae5f5cd1b3b..6fe25d7e48d 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
@@ -17,9 +17,9 @@
 #include "split/kernel_holdout_emission_blurring_pathtermination_ao.h"
 
 __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
-        ccl_global char *shader_data,          /* Required throughout the kernel except probabilistic path termination and AO */
+        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 */
         ccl_global float3 *throughput_coop,    /* Required for handling holdout material and AO */
@@ -75,9 +75,9 @@ __kernel void kernel_ocl_path_trace_holdout_emission_blurring_pathtermination_ao
 	if(ray_index != QUEUE_EMPTY_SLOT) {
 #endif
 		kernel_holdout_emission_blurring_pathtermination_ao(
-		        globals,
+		        (KernelGlobals *)kg,
 		        data,
-		        shader_data,
+		        (ShaderData *)sd,
 		        per_sample_output_buffers,
 		        rng_coop,
 		        throughput_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
index 1bc7808d834..fb327a32247 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_lamp_emission.cl
@@ -17,9 +17,9 @@
 #include "split/kernel_lamp_emission.h"
 
 __kernel void kernel_ocl_path_trace_lamp_emission(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
-        ccl_global char *shader_data,          /* Required for lamp emission */
+        ccl_global char *sd,                   /* Required for lamp emission */
         ccl_global float3 *throughput_coop,    /* Required for lamp emission */
         PathRadiance *PathRadiance_coop,       /* Required for lamp emission */
         ccl_global Ray *Ray_coop,              /* Required for lamp emission */
@@ -68,9 +68,9 @@ __kernel void kernel_ocl_path_trace_lamp_emission(
 		}
 	}
 
-	kernel_lamp_emission(globals,
+	kernel_lamp_emission((KenrelGLobals *)kg,
 	                     data,
-	                     shader_data,
+	                     (ShaderData *)sd,
 	                     throughput_coop,
 	                     PathRadiance_coop,
 	                     Ray_coop,
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 dcf4db40411..8896fe739d1 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_next_iteration_setup.cl
@@ -17,9 +17,9 @@
 #include "split/kernel_next_iteration_setup.h"
 
 __kernel void kernel_ocl_path_trace_next_iteration_setup(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
-        ccl_global char *shader_data,         /* Required for setting up ray for next iteration */
+        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 */
         PathRadiance *PathRadiance_coop,      /* Required for setting up ray for next iteration */
@@ -83,9 +83,9 @@ __kernel void kernel_ocl_path_trace_next_iteration_setup(
 #ifndef __COMPUTE_DEVICE_GPU__
 	if(ray_index != QUEUE_EMPTY_SLOT) {
 #endif
-		enqueue_flag = kernel_next_iteration_setup(globals,
+		enqueue_flag = kernel_next_iteration_setup((KernelGlobals *)kg,
 		                                           data,
-		                                           shader_data,
+		                                           (ShaderData *)sd,
 		                                           rng_coop,
 		                                           throughput_coop,
 		                                           PathRadiance_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
index e5fad7bce50..3750a0d0062 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_scene_intersect.cl
@@ -17,7 +17,7 @@
 #include "split/kernel_scene_intersect.h"
 
 __kernel void kernel_ocl_path_trace_scene_intersect(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
         ccl_global uint *rng_coop,
         ccl_global Ray *Ray_coop,              /* Required for scene_intersect */
@@ -65,7 +65,7 @@ __kernel void kernel_ocl_path_trace_scene_intersect(
 		}
 	}
 
-	kernel_scene_intersect(globals,
+	kernel_scene_intersect((KernelGlobals *)kg,
 	                       data,
 	                       rng_coop,
 	                       Ray_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
index b9f616e6bdf..2fa5496a84b 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shader_eval.cl
@@ -17,9 +17,9 @@
 #include "split/kernel_shader_eval.h"
 
 __kernel void kernel_ocl_path_trace_shader_eval(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
-        ccl_global char *shader_data,          /* Output ShaderData structure to be filled */
+        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 */
         ccl_global PathState *PathState_coop,  /* Required for all functions in this kernel */
@@ -57,9 +57,9 @@ __kernel void kernel_ocl_path_trace_shader_eval(
 	                        Queue_index);
 
 	/* Continue on with shader evaluation. */
-	kernel_shader_eval(globals,
+	kernel_shader_eval((KernelGlobals *)kg,
 	                   data,
-	                   shader_data,
+	                   (ShaderData *)sd,
 	                   rng_coop,
 	                   Ray_coop,
 	                   PathState_coop,
diff --git a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
index 03886c0a030..b364c769840 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel_shadow_blocked.cl
@@ -17,9 +17,9 @@
 #include "split/kernel_shadow_blocked.h"
 
 __kernel void kernel_ocl_path_trace_shadow_blocked(
-        ccl_global char *globals,
+        ccl_global char *kg,
         ccl_constant KernelData *data,
-        ccl_global char *shader_shadow,        /* Required for shadow blocked */
+        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 */
         ccl_global Ray *LightRay_ao_coop,      /* Required for AO's shadow blocked */
@@ -68,9 +68,9 @@ __kernel void kernel_ocl_path_trace_shadow_blocked(
 	if(ray_index == QUEUE_EMPTY_SLOT)
 		return;
 
-	kernel_shadow_blocked(globals,
+	kernel_shadow_blocked((KernelGlobals *)kg,
 	                      data,
-	                      shader_shadow,
+	                      (ShaderData *)sd_shadow,
 	                      PathState_coop,
 	                      LightRay_dl_coop,
 	                      LightRay_ao_coop,