diff options
author | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2016-07-30 16:18:21 +0300 |
---|---|---|
committer | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2016-07-30 19:25:52 +0300 |
commit | 6dc72b3ce6943bb444a6d6a4029cc8ce5b5bdf45 (patch) | |
tree | bb6b09e1e12d6f06b5ff172a43cb5ae38d7bf43d /intern/cycles/kernel/kernel_types.h | |
parent | c937a42c6171dc42c7ce477f2b794ce04f02763e (diff) |
Cycles OpenCL: detect incorrect usage of SOA members in the split kernel.
Diffstat (limited to 'intern/cycles/kernel/kernel_types.h')
-rw-r--r-- | intern/cycles/kernel/kernel_types.h | 63 |
1 files changed, 33 insertions, 30 deletions
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h index fd2e70b6d71..0be381ba800 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; |