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;
_______________________________________________
Bf-blender-cvs mailing list
[email protected]
https://lists.blender.org/mailman/listinfo/bf-blender-cvs
