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, 87 insertions, 10 deletions

diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index da2416bc09b..49d8859c7de 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -718,22 +718,99 @@ enum ShaderDataFlag {
 struct KernelGlobals;
 
 #ifdef __SPLIT_KERNEL__
-#define SD_VAR(type, what) ccl_global type *what;
-#define SD_CLOSURE_VAR(type, what, max_closure) type *what;
-#define TIDX (get_global_id(1) * get_global_size(0) + get_global_id(0))
-#define ccl_fetch(s, t) (s->t[TIDX])
-#define ccl_fetch_array(s, t, index) (&s->t[TIDX * MAX_CLOSURE + index])
+#  define SD_THREAD (get_global_id(1) * get_global_size(0) + get_global_id(0))
+#  if defined(__SPLIT_KERNEL_AOS__)
+     /* ShaderData is stored as an Array-of-Structures */
+#    define ccl_fetch(s, t) (s[SD_THREAD].t)
+#    define ccl_fetch_array(s, t, index) (&s[SD_THREAD].t[index])
+#  else
+     /* ShaderData is stored as an Structure-of-Arrays */
+#    define SD_GLOBAL_SIZE (get_global_size(0) * get_global_size(1))
+#    define SD_FIELD_SIZE(t) sizeof(((struct ShaderData*)0)->t)
+#    define SD_OFFSETOF(t) ((char*)(&((struct ShaderData*)0)->t) - (char*)0)
+#    define ccl_fetch(s, t) (((ShaderData*)((ccl_addr_space char*)s + SD_GLOBAL_SIZE * SD_OFFSETOF(t) +  SD_FIELD_SIZE(t) * SD_THREAD - SD_OFFSETOF(t)))->t)
+#    define ccl_fetch_array(s, t, index) (&ccl_fetch(s, t)[index])
+#  endif
 #else
-#define SD_VAR(type, what) type what;
-#define SD_CLOSURE_VAR(type, what, max_closure) type what[max_closure];
-#define ccl_fetch(s, t) (s->t)
-#define ccl_fetch_array(s, t, index) (&s->t[index])
+#  define ccl_fetch(s, t) (s->t)
+#  define ccl_fetch_array(s, t, index) (&s->t[index])
 #endif
 
 typedef ccl_addr_space struct ShaderData {
+	/* position */
+	float3 P;
+	/* smooth normal for shading */
+	float3 N;
+	/* true geometric normal */
+	float3 Ng;
+	/* view/incoming direction */
+	float3 I;
+	/* shader id */
+	int shader;
+	/* booleans describing shader, see ShaderDataFlag */
+	int flag;
+
+	/* primitive id if there is one, ~0 otherwise */
+	int prim;
+
+	/* combined type and curve segment for hair */
+	int type;
+
+	/* parametric coordinates
+	* - barycentric weights for triangles */
+	float u;
+	float v;
+	/* object id if there is one, ~0 otherwise */
+	int object;
+
+	/* motion blur sample time */
+	float time;
+
+	/* length of the ray being shaded */
+	float ray_length;
 
-#include "kernel_shaderdata_vars.h"
+	/* ray bounce depth */
+	int ray_depth;
 
+	/* ray transparent depth */
+	int transparent_depth;
+
+#ifdef __RAY_DIFFERENTIALS__
+	/* differential of P. these are orthogonal to Ng, not N */
+	differential3 dP;
+	/* differential of I */
+	differential3 dI;
+	/* differential of u, v */
+	differential du;
+	differential dv;
+#endif
+#ifdef __DPDU__
+	/* differential of P w.r.t. parametric coordinates. note that dPdu is
+	* not readily suitable as a tangent for shading on triangles. */
+	float3 dPdu;
+	float3 dPdv;
+#endif
+
+#ifdef __OBJECT_MOTION__
+	/* object <-> world space transformations, cached to avoid
+	* re-interpolating them constantly for shading */
+	Transform ob_tfm;
+	Transform ob_itfm;
+#endif
+
+	/* Closure data, we store a fixed array of closures */
+	struct ShaderClosure closure[MAX_CLOSURE];
+	int num_closure;
+	float randb_closure;
+
+	/* ray start position, only set for backgrounds */
+	float3 ray_P;
+	differential3 ray_dP;
+
+#ifdef __OSL__
+	struct KernelGlobals * osl_globals;
+	struct PathState *osl_path_state;
+#endif
 } ShaderData;
 
 /* Path State */