blender
/
blender
130
397
Fork 571
Code Issues 6.2k Pull Requests 756 Projects 19 Wiki Activity

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

This commit is contained in:
Brecht Van Lommel 2016-07-30 15:18:21 +02:00
parent c937a42c61
commit 6dc72b3ce6
1 changed files with 33 additions and 30 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

63
intern/cycles/kernel/kernel_types.h
View File

@ -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;