Cycles: Cleanup of OpenCL split kernel routines

The idea is to switch from allocating separate buffers for shader data's
structure of arrays to allocating one huge memory block and do some index
trickery to make it accessed as SOA.

This saves quite reasonable amount of lines of code in device_opencl and
also makes it possible to get rid of special declaration of ShaderData
structure.

As a side effect it also makes it easier to experiment with SOA vs. AOS
for split kernel.

Works fine here on NVidia GTX580, Intel CPU amd AMD Fiji cards.

Reviewers: #cycles, brecht, juicyfruit, dingto

Differential Revision: https://developer.blender.org/D1593

1 files changed, 0 insertions, 166 deletions

diff --git a/intern/cycles/kernel/split/kernel_data_init.h b/intern/cycles/kernel/split/kernel_data_init.h
index 6993a78a789..9891391a3a3 100644
--- a/intern/cycles/kernel/split/kernel_data_init.h
+++ b/intern/cycles/kernel/split/kernel_data_init.h
@@ -52,90 +52,7 @@
  */
 ccl_device void kernel_data_init(
         KernelGlobals *kg,
-        ShaderData *sd,
         ShaderData *sd_DL_shadow,
-
-        ccl_global float3 *P_sd,
-        ccl_global float3 *P_sd_DL_shadow,
-
-        ccl_global float3 *N_sd,
-        ccl_global float3 *N_sd_DL_shadow,
-
-        ccl_global float3 *Ng_sd,
-        ccl_global float3 *Ng_sd_DL_shadow,
-
-        ccl_global float3 *I_sd,
-        ccl_global float3 *I_sd_DL_shadow,
-
-        ccl_global int *shader_sd,
-        ccl_global int *shader_sd_DL_shadow,
-
-        ccl_global int *flag_sd,
-        ccl_global int *flag_sd_DL_shadow,
-
-        ccl_global int *prim_sd,
-        ccl_global int *prim_sd_DL_shadow,
-
-        ccl_global int *type_sd,
-        ccl_global int *type_sd_DL_shadow,
-
-        ccl_global float *u_sd,
-        ccl_global float *u_sd_DL_shadow,
-
-        ccl_global float *v_sd,
-        ccl_global float *v_sd_DL_shadow,
-
-        ccl_global int *object_sd,
-        ccl_global int *object_sd_DL_shadow,
-
-        ccl_global float *time_sd,
-        ccl_global float *time_sd_DL_shadow,
-
-        ccl_global float *ray_length_sd,
-        ccl_global float *ray_length_sd_DL_shadow,
-
-        /* Ray differentials. */
-        ccl_global differential3 *dP_sd,
-        ccl_global differential3 *dP_sd_DL_shadow,
-
-        ccl_global differential3 *dI_sd,
-        ccl_global differential3 *dI_sd_DL_shadow,
-
-        ccl_global differential *du_sd,
-        ccl_global differential *du_sd_DL_shadow,
-
-        ccl_global differential *dv_sd,
-        ccl_global differential *dv_sd_DL_shadow,
-
-        /* Dp/Du */
-        ccl_global float3 *dPdu_sd,
-        ccl_global float3 *dPdu_sd_DL_shadow,
-
-        ccl_global float3 *dPdv_sd,
-        ccl_global float3 *dPdv_sd_DL_shadow,
-
-        /* Object motion. */
-        ccl_global Transform *ob_tfm_sd,
-        ccl_global Transform *ob_tfm_sd_DL_shadow,
-
-        ccl_global Transform *ob_itfm_sd,
-        ccl_global Transform *ob_itfm_sd_DL_shadow,
-
-        ShaderClosure *closure_sd,
-        ShaderClosure *closure_sd_DL_shadow,
-
-        ccl_global int *num_closure_sd,
-        ccl_global int *num_closure_sd_DL_shadow,
-
-        ccl_global float *randb_closure_sd,
-        ccl_global float *randb_closure_sd_DL_shadow,
-
-        ccl_global float3 *ray_P_sd,
-        ccl_global float3 *ray_P_sd_DL_shadow,
-
-        ccl_global differential3 *ray_dP_sd,
-        ccl_global differential3 *ray_dP_sd_DL_shadow,
-
         ccl_constant KernelData *data,
         ccl_global float *per_sample_output_buffers,
         ccl_global uint *rng_state,
@@ -177,89 +94,6 @@ ccl_device void kernel_data_init(
 	kg->name = name;
 #include "../kernel_textures.h"
 
-	sd->P = P_sd;
-	sd_DL_shadow->P = P_sd_DL_shadow;
-
-	sd->N = N_sd;
-	sd_DL_shadow->N = N_sd_DL_shadow;
-
-	sd->Ng = Ng_sd;
-	sd_DL_shadow->Ng = Ng_sd_DL_shadow;
-
-	sd->I = I_sd;
-	sd_DL_shadow->I = I_sd_DL_shadow;
-
-	sd->shader = shader_sd;
-	sd_DL_shadow->shader = shader_sd_DL_shadow;
-
-	sd->flag = flag_sd;
-	sd_DL_shadow->flag = flag_sd_DL_shadow;
-
-	sd->prim = prim_sd;
-	sd_DL_shadow->prim = prim_sd_DL_shadow;
-
-	sd->type = type_sd;
-	sd_DL_shadow->type = type_sd_DL_shadow;
-
-	sd->u = u_sd;
-	sd_DL_shadow->u = u_sd_DL_shadow;
-
-	sd->v = v_sd;
-	sd_DL_shadow->v = v_sd_DL_shadow;
-
-	sd->object = object_sd;
-	sd_DL_shadow->object = object_sd_DL_shadow;
-
-	sd->time = time_sd;
-	sd_DL_shadow->time = time_sd_DL_shadow;
-
-	sd->ray_length = ray_length_sd;
-	sd_DL_shadow->ray_length = ray_length_sd_DL_shadow;
-
-#ifdef __RAY_DIFFERENTIALS__
-	sd->dP = dP_sd;
-	sd_DL_shadow->dP = dP_sd_DL_shadow;
-
-	sd->dI = dI_sd;
-	sd_DL_shadow->dI = dI_sd_DL_shadow;
-
-	sd->du = du_sd;
-	sd_DL_shadow->du = du_sd_DL_shadow;
-
-	sd->dv = dv_sd;
-	sd_DL_shadow->dv = dv_sd_DL_shadow;
-#ifdef __DPDU__
-	sd->dPdu = dPdu_sd;
-	sd_DL_shadow->dPdu = dPdu_sd_DL_shadow;
-
-	sd->dPdv = dPdv_sd;
-	sd_DL_shadow->dPdv = dPdv_sd_DL_shadow;
-#endif
-#endif
-
-#ifdef __OBJECT_MOTION__
-	sd->ob_tfm = ob_tfm_sd;
-	sd_DL_shadow->ob_tfm = ob_tfm_sd_DL_shadow;
-
-	sd->ob_itfm = ob_itfm_sd;
-	sd_DL_shadow->ob_itfm = ob_itfm_sd_DL_shadow;
-#endif
-
-	sd->closure = closure_sd;
-	sd_DL_shadow->closure = closure_sd_DL_shadow;
-
-	sd->num_closure = num_closure_sd;
-	sd_DL_shadow->num_closure = num_closure_sd_DL_shadow;
-
-	sd->randb_closure = randb_closure_sd;
-	sd_DL_shadow->randb_closure = randb_closure_sd_DL_shadow;
-
-	sd->ray_P = ray_P_sd;
-	sd_DL_shadow->ray_P = ray_P_sd_DL_shadow;
-
-	sd->ray_dP = ray_dP_sd;
-	sd_DL_shadow->ray_dP = ray_dP_sd_DL_shadow;
-
 	int thread_index = get_global_id(1) * get_global_size(0) + get_global_id(0);
 
 #ifdef __WORK_STEALING__