diff options
author | Pascal Schoen <pascal_schoen@gmx.net> | 2016-08-03 12:42:02 +0300 |
---|---|---|
committer | Pascal Schoen <pascal_schoen@gmx.net> | 2016-08-03 12:42:02 +0300 |
commit | 81f6c06b1f53180bf32a5c11ac1fa64e2b6abf52 (patch) | |
tree | c7ad4920e48e0eb529e2064fd0d3813c29d5383b /intern/cycles/kernel/kernel_types.h | |
parent | ece5a08e0d6e51a83c223ea87346134216e5b34e (diff) | |
parent | 7065022f7aa23ba13d2999e1e40162a8f480af0e (diff) |
Merge branch 'master' into cycles_disney_brdf
Diffstat (limited to 'intern/cycles/kernel/kernel_types.h')
-rw-r--r-- | intern/cycles/kernel/kernel_types.h | 235 |
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 */ |