Merge branch 'master' into cycles_disney_brdf

1 files changed, 120 insertions, 115 deletions

diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index df5e5d803fa..18b5c35c768 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -121,9 +121,7 @@ CCL_NAMESPACE_BEGIN
 #    define __OBJECT_MOTION__
 #    define __HAIR__
 #    define __BAKING__
-#    ifdef __KERNEL_EXPERIMENTAL__
-#      define __TRANSPARENT_SHADOWS__
-#    endif
+#    define __TRANSPARENT_SHADOWS__
 #  endif  /* __KERNEL_OPENCL_AMD__ */
 
 #  ifdef __KERNEL_OPENCL_INTEL_CPU__
@@ -294,12 +292,14 @@ enum PathRayFlag {
 	PATH_RAY_CURVE = 512, /* visibility flag to define curve segments */
 	PATH_RAY_VOLUME_SCATTER = 1024, /* volume scattering */
 
-	/* note that these can use maximum 12 bits, the other are for layers */
-	PATH_RAY_ALL_VISIBILITY = (1|2|4|8|16|32|64|128|256|512|1024),
+	/* Special flag to tag unaligned BVH nodes. */
+	PATH_RAY_NODE_UNALIGNED = 2048,
+
+	PATH_RAY_ALL_VISIBILITY = (1|2|4|8|16|32|64|128|256|512|1024|2048),
 
-	PATH_RAY_MIS_SKIP = 2048,
-	PATH_RAY_DIFFUSE_ANCESTOR = 4096,
-	PATH_RAY_SINGLE_PASS_DONE = 8192,
+	PATH_RAY_MIS_SKIP = 4096,
+	PATH_RAY_DIFFUSE_ANCESTOR = 8192,
+	PATH_RAY_SINGLE_PASS_DONE = 16384,
 };
 
 /* Closure Label */
@@ -387,12 +387,13 @@ typedef enum BakePassFilterCombos {
 	BAKE_FILTER_SUBSURFACE_INDIRECT = (BAKE_FILTER_INDIRECT | BAKE_FILTER_SUBSURFACE),
 } BakePassFilterCombos;
 
-#ifdef __PASSES__
-
 typedef ccl_addr_space struct PathRadiance {
+#ifdef __PASSES__
 	int use_light_pass;
+#endif
 
 	float3 emission;
+#ifdef __PASSES__
 	float3 background;
 	float3 ao;
 
@@ -426,25 +427,23 @@ typedef ccl_addr_space struct PathRadiance {
 
 	float4 shadow;
 	float mist;
+#endif
 } PathRadiance;
 
 typedef struct BsdfEval {
+#ifdef __PASSES__
 	int use_light_pass;
+#endif
 
 	float3 diffuse;
+#ifdef __PASSES__
 	float3 glossy;
 	float3 transmission;
 	float3 transparent;
 	float3 subsurface;
 	float3 scatter;
-} BsdfEval;
-
-#else
-
-typedef ccl_addr_space float3 PathRadiance;
-typedef float3 BsdfEval;
-
 #endif
+} BsdfEval;
 
 /* Shader Flag */
 
@@ -574,8 +573,13 @@ typedef enum PrimitiveType {
 
 /* Attributes */
 
-#define ATTR_PRIM_TYPES		2
-#define ATTR_PRIM_CURVE		1
+typedef enum AttributePrimitive {
+	ATTR_PRIM_TRIANGLE = 0,
+	ATTR_PRIM_CURVE,
+	ATTR_PRIM_SUBD,
+
+	ATTR_PRIM_TYPES
+} AttributePrimitive;
 
 typedef enum AttributeElement {
 	ATTR_ELEMENT_NONE,
@@ -632,37 +636,30 @@ typedef enum AttributeStandard {
 #  define MAX_CLOSURE 1
 #endif
 
-/* This struct is to be 16 bytes aligned, we also keep some extra precautions:
- * - All the float3 members are in the beginning of the struct, so compiler
- *   does not put own padding trying to align this members.
- * - We make sure OSL pointer is also 16 bytes aligned.
- */
+/* This struct is the base class for all closures. The common members are
+ * duplicated in all derived classes since we don't have C++ in the kernel
+ * yet, and because it lets us lay out the members to minimize padding. The
+ * weight member is located at the beginning of the struct for this reason.
+ *
+ * ShaderClosure has a fixed size, and any extra space must be allocated
+ * with closure_alloc_extra().
+ *
+ * float3 is 12 bytes on CUDA and 16 bytes on CPU/OpenCL, we set the data
+ * size to ensure ShaderClosure is 80 bytes total everywhere. */
+
+#define SHADER_CLOSURE_BASE \
+	float3 weight; \
+	ClosureType type; \
+	float sample_weight \
+
 typedef ccl_addr_space struct ShaderClosure {
-	float3 weight;
-	float3 N;
-	float3 T;
-
-	float3 color0;
-
-	ClosureType type;
-	float sample_weight;
-	float data0;
-	float data1;
-	float data2;
-	float data3;
-	float data4;
-
-	/* Following fields could be used to store pre-calculated
-	 * values by various BSDF closures for more effective sampling
-	 * and evaluation.
-	 */
-	float custom1;
-	float custom2;
-	float custom3;
-	float3 custom_color0;
+	SHADER_CLOSURE_BASE;
 
-#ifdef __OSL__
-	void *prim, *pad4;
+	/* pad to 80 bytes, data types are aligned to own size */
+#ifdef __KERNEL_CUDA__
+	float data[15];
+#else
+	float data[14];
 #endif
 } ShaderClosure;
 
@@ -687,31 +684,33 @@ typedef enum ShaderContext {
 
 enum ShaderDataFlag {
 	/* runtime flags */
-	SD_BACKFACING     = (1 << 0),   /* backside of surface? */
-	SD_EMISSION       = (1 << 1),   /* have emissive closure? */
-	SD_BSDF           = (1 << 2),   /* have bsdf closure? */
-	SD_BSDF_HAS_EVAL  = (1 << 3),   /* have non-singular bsdf closure? */
-	SD_BSSRDF         = (1 << 4),   /* have bssrdf */
-	SD_HOLDOUT        = (1 << 5),   /* have holdout closure? */
-	SD_ABSORPTION     = (1 << 6),   /* have volume absorption closure? */
-	SD_SCATTER        = (1 << 7),   /* have volume phase closure? */
-	SD_AO             = (1 << 8),   /* have ao closure? */
-	SD_TRANSPARENT    = (1 << 9),  /* have transparent closure? */
+	SD_BACKFACING      = (1 << 0),   /* backside of surface? */
+	SD_EMISSION        = (1 << 1),   /* have emissive closure? */
+	SD_BSDF            = (1 << 2),   /* have bsdf closure? */
+	SD_BSDF_HAS_EVAL   = (1 << 3),   /* have non-singular bsdf closure? */
+	SD_BSSRDF          = (1 << 4),   /* have bssrdf */
+	SD_HOLDOUT         = (1 << 5),   /* have holdout closure? */
+	SD_ABSORPTION      = (1 << 6),   /* have volume absorption closure? */
+	SD_SCATTER         = (1 << 7),   /* have volume phase closure? */
+	SD_AO              = (1 << 8),   /* have ao closure? */
+	SD_TRANSPARENT     = (1 << 9),  /* have transparent closure? */
+	SD_BSDF_NEEDS_LCG  = (1 << 10),
 
 	SD_CLOSURE_FLAGS = (SD_EMISSION|SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSSRDF|
-	                    SD_HOLDOUT|SD_ABSORPTION|SD_SCATTER|SD_AO),
+	                    SD_HOLDOUT|SD_ABSORPTION|SD_SCATTER|SD_AO|
+	                    SD_BSDF_NEEDS_LCG),
 
 	/* shader flags */
-	SD_USE_MIS                = (1 << 10),  /* direct light sample */
-	SD_HAS_TRANSPARENT_SHADOW = (1 << 11),  /* has transparent shadow */
-	SD_HAS_VOLUME             = (1 << 12),  /* has volume shader */
-	SD_HAS_ONLY_VOLUME        = (1 << 13),  /* has only volume shader, no surface */
-	SD_HETEROGENEOUS_VOLUME   = (1 << 14),  /* has heterogeneous volume */
-	SD_HAS_BSSRDF_BUMP        = (1 << 15),  /* bssrdf normal uses bump */
-	SD_VOLUME_EQUIANGULAR     = (1 << 16),  /* use equiangular sampling */
-	SD_VOLUME_MIS             = (1 << 17),  /* use multiple importance sampling */
-	SD_VOLUME_CUBIC           = (1 << 18),  /* use cubic interpolation for voxels */
-	SD_HAS_BUMP               = (1 << 19),  /* has data connected to the displacement input */
+	SD_USE_MIS                = (1 << 12),  /* direct light sample */
+	SD_HAS_TRANSPARENT_SHADOW = (1 << 13),  /* has transparent shadow */
+	SD_HAS_VOLUME             = (1 << 14),  /* has volume shader */
+	SD_HAS_ONLY_VOLUME        = (1 << 15),  /* has only volume shader, no surface */
+	SD_HETEROGENEOUS_VOLUME   = (1 << 16),  /* has heterogeneous volume */
+	SD_HAS_BSSRDF_BUMP        = (1 << 17),  /* bssrdf normal uses bump */
+	SD_VOLUME_EQUIANGULAR     = (1 << 18),  /* use equiangular sampling */
+	SD_VOLUME_MIS             = (1 << 19),  /* use multiple importance sampling */
+	SD_VOLUME_CUBIC           = (1 << 20),  /* use cubic interpolation for voxels */
+	SD_HAS_BUMP               = (1 << 21),  /* has data connected to the displacement input */
 
 	SD_SHADER_FLAGS = (SD_USE_MIS|SD_HAS_TRANSPARENT_SHADOW|SD_HAS_VOLUME|
 	                   SD_HAS_ONLY_VOLUME|SD_HETEROGENEOUS_VOLUME|
@@ -719,104 +718,110 @@ enum ShaderDataFlag {
 	                   SD_VOLUME_CUBIC|SD_HAS_BUMP),
 
 	/* object flags */
-	SD_HOLDOUT_MASK             = (1 << 20),  /* holdout for camera rays */
-	SD_OBJECT_MOTION            = (1 << 21),  /* has object motion blur */
-	SD_TRANSFORM_APPLIED        = (1 << 22),  /* vertices have transform applied */
-	SD_NEGATIVE_SCALE_APPLIED   = (1 << 23),  /* vertices have negative scale applied */
-	SD_OBJECT_HAS_VOLUME        = (1 << 24),  /* object has a volume shader */
-	SD_OBJECT_INTERSECTS_VOLUME = (1 << 25),  /* object intersects AABB of an object with volume shader */
-	SD_OBJECT_HAS_VERTEX_MOTION = (1 << 26),  /* has position for motion vertices */
+	SD_HOLDOUT_MASK             = (1 << 22),  /* holdout for camera rays */
+	SD_OBJECT_MOTION            = (1 << 23),  /* has object motion blur */
+	SD_TRANSFORM_APPLIED        = (1 << 24),  /* vertices have transform applied */
+	SD_NEGATIVE_SCALE_APPLIED   = (1 << 25),  /* vertices have negative scale applied */
+	SD_OBJECT_HAS_VOLUME        = (1 << 26),  /* object has a volume shader */
+	SD_OBJECT_INTERSECTS_VOLUME = (1 << 27),  /* object intersects AABB of an object with volume shader */
+	SD_OBJECT_HAS_VERTEX_MOTION = (1 << 28),  /* has position for motion vertices */
 
 	SD_OBJECT_FLAGS = (SD_HOLDOUT_MASK|SD_OBJECT_MOTION|SD_TRANSFORM_APPLIED|
 	                   SD_NEGATIVE_SCALE_APPLIED|SD_OBJECT_HAS_VOLUME|
 	                   SD_OBJECT_INTERSECTS_VOLUME)
 };
 
-struct KernelGlobals;
-
 #ifdef __SPLIT_KERNEL__
 #  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;
+	 * - barycentric weights for triangles */
+	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;
+	 * not readily suitable as a tangent for shading on triangles. */
+	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;
+	 * re-interpolating them constantly for shading */
+	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(int, num_closure_extra);
+	ccl_soa_member(float, randb_closure);
+	ccl_soa_member(float3, svm_closure_weight);
+
+	/* LCG state for closures that require additional random numbers. */
+	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;
@@ -1173,11 +1178,11 @@ typedef ccl_addr_space struct DebugData {
 #define QUEUE_EMPTY_SLOT -1
 
 /*
-* Queue 1 - Active rays
-* Queue 2 - Background queue
-* Queue 3 - Shadow ray cast kernel - AO
-* Queeu 4 - Shadow ray cast kernel - direct lighting
-*/
+ * Queue 1 - Active rays
+ * Queue 2 - Background queue
+ * Queue 3 - Shadow ray cast kernel - AO
+ * Queeu 4 - Shadow ray cast kernel - direct lighting
+ */
 #define NUM_QUEUES 4
 
 /* Queue names */