[Bf-blender-cvs] [6dc72b3] master: Cycles OpenCL: detect incorrect usage of SOA members in the split kernel.

[Brecht Van Lommel]

Commit: 6dc72b3ce6943bb444a6d6a4029cc8ce5b5bdf45
Author: Brecht Van Lommel
Date:   Sat Jul 30 15:18:21 2016 +0200
Branches: master
https://developer.blender.org/rB6dc72b3ce6943bb444a6d6a4029cc8ce5b5bdf45

Cycles OpenCL: detect incorrect usage of SOA members in the split kernel.

===================================================================

M	intern/cycles/kernel/kernel_types.h

===================================================================

diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index fd2e70b..0be381b 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -738,88 +738,91 @@ enum ShaderDataFlag {
 #  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])
+#    define ccl_soa_member(type, name) type soa_##name;
+#    define ccl_fetch(s, t) (s[SD_THREAD].soa_##t)
+#    define ccl_fetch_array(s, t, index) (&s[SD_THREAD].soa_##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_soa_member(type, name) type soa_##name;
+#    define ccl_fetch(s, t) (((ShaderData*)((ccl_addr_space char*)s + SD_GLOBAL_SIZE * SD_OFFSETOF(soa_##t) +  SD_FIELD_SIZE(soa_##t) * SD_THREAD - SD_OFFSETOF(soa_##t)))->soa_##t)
 #    define ccl_fetch_array(s, t, index) (&ccl_fetch(s, t)[index])
 #  endif
 #else
+#  define ccl_soa_member(type, name) type name;
 #  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;
+	ccl_soa_member(float3, P);
 	/* smooth normal for shading */
-	float3 N;
+	ccl_soa_member(float3, N);
 	/* true geometric normal */
-	float3 Ng;
+	ccl_soa_member(float3, Ng);
 	/* view/incoming direction */
-	float3 I;
+	ccl_soa_member(float3, I);
 	/* shader id */
-	int shader;
+	ccl_soa_member(int, shader);
 	/* booleans describing shader, see ShaderDataFlag */
-	int flag;
+	ccl_soa_member(int, flag);
 
 	/* primitive id if there is one, ~0 otherwise */
-	int prim;
+	ccl_soa_member(int, prim);
 
 	/* combined type and curve segment for hair */
-	int type;
+	ccl_soa_member(int, type);
 
 	/* parametric coordinates
 	 * - barycentric weights for triangles */
-	float u;
-	float v;
+	ccl_soa_member(float, u);
+	ccl_soa_member(float, v);
 	/* object id if there is one, ~0 otherwise */
-	int object;
+	ccl_soa_member(int, object);
 
 	/* motion blur sample time */
-	float time;
+	ccl_soa_member(float, time);
 
 	/* length of the ray being shaded */
-	float ray_length;
+	ccl_soa_member(float, ray_length);
 
 #ifdef __RAY_DIFFERENTIALS__
 	/* differential of P. these are orthogonal to Ng, not N */
-	differential3 dP;
+	ccl_soa_member(differential3, dP);
 	/* differential of I */
-	differential3 dI;
+	ccl_soa_member(differential3, dI);
 	/* differential of u, v */
-	differential du;
-	differential dv;
+	ccl_soa_member(differential, du);
+	ccl_soa_member(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;
+	ccl_soa_member(float3, dPdu);
+	ccl_soa_member(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;
+	ccl_soa_member(Transform, ob_tfm);
+	ccl_soa_member(Transform, ob_itfm);
 #endif
 
 	/* Closure data, we store a fixed array of closures */
-	struct ShaderClosure closure[MAX_CLOSURE];
-	int num_closure;
-	float randb_closure;
+	ccl_soa_member(struct ShaderClosure, closure[MAX_CLOSURE]);
+	ccl_soa_member(int, num_closure);
+	ccl_soa_member(float, randb_closure);
 
 	/* LCG state for closures that require additional random numbers. */
-	uint lcg_state;
+	ccl_soa_member(uint, lcg_state);
 
 	/* ray start position, only set for backgrounds */
-	float3 ray_P;
-	differential3 ray_dP;
+	ccl_soa_member(float3, ray_P);
+	ccl_soa_member(differential3, ray_dP);
 
 #ifdef __OSL__
 	struct KernelGlobals * osl_globals;