Cycles: Workaround for AMD compiler crashing building the split kernel
[blender-staging.git] / intern / cycles / kernel / kernel_types.h
1 /*
2  * Copyright 2011-2013 Blender Foundation
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16
17 #ifndef __KERNEL_TYPES_H__
18 #define __KERNEL_TYPES_H__
19
20 #include "kernel_math.h"
21 #include "svm/svm_types.h"
22
23 #ifndef __KERNEL_GPU__
24 #define __KERNEL_CPU__
25 #endif
26
27 /* TODO(sergey): This is only to make it possible to include this header
28  * from outside of the kernel. but this could be done somewhat cleaner?
29  */
30 #ifndef ccl_addr_space
31 #define ccl_addr_space
32 #endif
33
34 CCL_NAMESPACE_BEGIN
35
36 /* constants */
37 #define OBJECT_SIZE             11
38 #define OBJECT_VECTOR_SIZE      6
39 #define LIGHT_SIZE                      5
40 #define FILTER_TABLE_SIZE       256
41 #define RAMP_TABLE_SIZE         256
42 #define PARTICLE_SIZE           5
43 #define TIME_INVALID            FLT_MAX
44
45 #define BSSRDF_MIN_RADIUS                       1e-8f
46 #define BSSRDF_MAX_HITS                         4
47
48 #define BECKMANN_TABLE_SIZE             256
49
50 #define TEX_NUM_FLOAT_IMAGES    5
51
52 #define SHADER_NONE                             (~0)
53 #define OBJECT_NONE                             (~0)
54 #define PRIM_NONE                               (~0)
55 #define LAMP_NONE                               (~0)
56
57 #define VOLUME_STACK_SIZE               16
58
59 /* device capabilities */
60 #ifdef __KERNEL_CPU__
61 #ifdef __KERNEL_SSE2__
62 #  define __QBVH__
63 #endif
64 #define __KERNEL_SHADING__
65 #define __KERNEL_ADV_SHADING__
66 #define __BRANCHED_PATH__
67 #ifdef WITH_OSL
68 #define __OSL__
69 #endif
70 #define __SUBSURFACE__
71 #define __CMJ__
72 #define __VOLUME__
73 #define __VOLUME_DECOUPLED__
74 #define __VOLUME_SCATTER__
75 #define __SHADOW_RECORD_ALL__
76 #define __VOLUME_RECORD_ALL__
77 #endif
78
79 #ifdef __KERNEL_CUDA__
80 #define __KERNEL_SHADING__
81 #define __KERNEL_ADV_SHADING__
82 #define __BRANCHED_PATH__
83 #define __VOLUME__
84 #define __VOLUME_SCATTER__
85
86 /* Experimental on GPU */
87 #ifdef __KERNEL_CUDA_EXPERIMENTAL__
88 #define __SUBSURFACE__
89 #define __CMJ__
90 #endif
91
92 #endif
93
94 #ifdef __KERNEL_OPENCL__
95
96 /* keep __KERNEL_ADV_SHADING__ in sync with opencl_kernel_use_advanced_shading! */
97
98 #ifdef __KERNEL_OPENCL_NVIDIA__
99 #define __KERNEL_SHADING__
100 /* TODO(sergey): Advanced shading code still requires work
101  * for split kernel.
102  */
103 #  ifndef __SPLIT_KERNEL__
104 #    define __KERNEL_ADV_SHADING__
105 #  else
106 #    define __MULTI_CLOSURE__
107 #    define __TRANSPARENT_SHADOWS__
108 #    define __PASSES__
109 #    define __BACKGROUND_MIS__
110 #    define __LAMP_MIS__
111 #    define __AO__
112 #  endif
113 #endif
114
115 #ifdef __KERNEL_OPENCL_APPLE__
116 #define __KERNEL_SHADING__
117 //#define __KERNEL_ADV_SHADING__
118 #endif
119
120 #ifdef __KERNEL_OPENCL_AMD__
121 #define __CL_USE_NATIVE__
122 #define __KERNEL_SHADING__
123 //__KERNEL_ADV_SHADING__
124 #define __MULTI_CLOSURE__
125 //#define __TRANSPARENT_SHADOWS__
126 #define __PASSES__
127 #define __BACKGROUND_MIS__
128 #define __LAMP_MIS__
129 #define __AO__
130 //#define __CAMERA_MOTION__
131 //#define __OBJECT_MOTION__
132 //#define __HAIR__
133 //end __KERNEL_ADV_SHADING__
134 #endif
135
136 #ifdef __KERNEL_OPENCL_INTEL_CPU__
137 #define __CL_USE_NATIVE__
138 #define __KERNEL_SHADING__
139 /* TODO(sergey): Advanced shading code still requires work
140  * for split kernel.
141  */
142 #  ifndef __SPLIT_KERNEL__
143 #    define __KERNEL_ADV_SHADING__
144 #  else
145 #    define __MULTI_CLOSURE__
146 #    define __TRANSPARENT_SHADOWS__
147 #    define __PASSES__
148 #    define __BACKGROUND_MIS__
149 #    define __LAMP_MIS__
150 #    define __AO__
151 #  endif
152 #endif
153
154 #endif // __KERNEL_OPENCL__
155
156 /* kernel features */
157 #define __SOBOL__
158 #define __INSTANCING__
159 #define __DPDU__
160 #define __UV__
161 #define __BACKGROUND__
162 #define __CAUSTICS_TRICKS__
163 #define __VISIBILITY_FLAG__
164 #define __RAY_DIFFERENTIALS__
165 #define __CAMERA_CLIPPING__
166 #define __INTERSECTION_REFINE__
167 #define __CLAMP_SAMPLE__
168
169 #ifdef __KERNEL_SHADING__
170 #define __SVM__
171 #define __EMISSION__
172 #define __TEXTURES__
173 #define __EXTRA_NODES__
174 #define __HOLDOUT__
175 #endif
176
177 #ifdef __KERNEL_ADV_SHADING__
178 #define __MULTI_CLOSURE__
179 #define __TRANSPARENT_SHADOWS__
180 #define __PASSES__
181 #define __BACKGROUND_MIS__
182 #define __LAMP_MIS__
183 #define __AO__
184 #define __CAMERA_MOTION__
185 #define __OBJECT_MOTION__
186 #define __HAIR__
187 #endif
188
189 #ifdef WITH_CYCLES_DEBUG
190 #  define __KERNEL_DEBUG__
191 #endif
192
193 /* Random Numbers */
194
195 typedef uint RNG;
196
197 /* Shader Evaluation */
198
199 typedef enum ShaderEvalType {
200         SHADER_EVAL_DISPLACE,
201         SHADER_EVAL_BACKGROUND,
202         /* bake types */
203         SHADER_EVAL_BAKE, /* no real shade, it's used in the code to
204                            * differentiate the type of shader eval from the above
205                            */
206         /* data passes */
207         SHADER_EVAL_NORMAL,
208         SHADER_EVAL_UV,
209         SHADER_EVAL_DIFFUSE_COLOR,
210         SHADER_EVAL_GLOSSY_COLOR,
211         SHADER_EVAL_TRANSMISSION_COLOR,
212         SHADER_EVAL_SUBSURFACE_COLOR,
213         SHADER_EVAL_EMISSION,
214
215         /* light passes */
216         SHADER_EVAL_AO,
217         SHADER_EVAL_COMBINED,
218         SHADER_EVAL_SHADOW,
219         SHADER_EVAL_DIFFUSE_DIRECT,
220         SHADER_EVAL_GLOSSY_DIRECT,
221         SHADER_EVAL_TRANSMISSION_DIRECT,
222         SHADER_EVAL_SUBSURFACE_DIRECT,
223         SHADER_EVAL_DIFFUSE_INDIRECT,
224         SHADER_EVAL_GLOSSY_INDIRECT,
225         SHADER_EVAL_TRANSMISSION_INDIRECT,
226         SHADER_EVAL_SUBSURFACE_INDIRECT,
227
228         /* extra */
229         SHADER_EVAL_ENVIRONMENT,
230 } ShaderEvalType;
231
232 /* Path Tracing
233  * note we need to keep the u/v pairs at even values */
234
235 enum PathTraceDimension {
236         PRNG_FILTER_U = 0,
237         PRNG_FILTER_V = 1,
238         PRNG_LENS_U = 2,
239         PRNG_LENS_V = 3,
240 #ifdef __CAMERA_MOTION__
241         PRNG_TIME = 4,
242         PRNG_UNUSED_0 = 5,
243         PRNG_UNUSED_1 = 6,      /* for some reason (6, 7) is a bad sobol pattern */
244         PRNG_UNUSED_2 = 7,  /* with a low number of samples (< 64) */
245 #endif
246         PRNG_BASE_NUM = 8,
247
248         PRNG_BSDF_U = 0,
249         PRNG_BSDF_V = 1,
250         PRNG_BSDF = 2,
251         PRNG_LIGHT = 3,
252         PRNG_LIGHT_U = 4,
253         PRNG_LIGHT_V = 5,
254         PRNG_UNUSED_3 = 6,
255         PRNG_TERMINATE = 7,
256
257 #ifdef __VOLUME__
258         PRNG_PHASE_U = 8,
259         PRNG_PHASE_V = 9,
260         PRNG_PHASE = 10,
261         PRNG_SCATTER_DISTANCE = 11,
262 #endif
263
264         PRNG_BOUNCE_NUM = 12,
265 };
266
267 enum SamplingPattern {
268         SAMPLING_PATTERN_SOBOL = 0,
269         SAMPLING_PATTERN_CMJ = 1
270 };
271
272 /* these flags values correspond to raytypes in osl.cpp, so keep them in sync!
273  *
274  * for ray visibility tests in BVH traversal, the upper 20 bits are used for
275  * layer visibility tests. */
276
277 enum PathRayFlag {
278         PATH_RAY_CAMERA = 1,
279         PATH_RAY_REFLECT = 2,
280         PATH_RAY_TRANSMIT = 4,
281         PATH_RAY_DIFFUSE = 8,
282         PATH_RAY_GLOSSY = 16,
283         PATH_RAY_SINGULAR = 32,
284         PATH_RAY_TRANSPARENT = 64,
285
286         PATH_RAY_SHADOW_OPAQUE = 128,
287         PATH_RAY_SHADOW_TRANSPARENT = 256,
288         PATH_RAY_SHADOW = (PATH_RAY_SHADOW_OPAQUE|PATH_RAY_SHADOW_TRANSPARENT),
289
290         PATH_RAY_CURVE = 512, /* visibility flag to define curve segments */
291         PATH_RAY_VOLUME_SCATTER = 1024, /* volume scattering */
292
293         /* note that these can use maximum 12 bits, the other are for layers */
294         PATH_RAY_ALL_VISIBILITY = (1|2|4|8|16|32|64|128|256|512|1024),
295
296         PATH_RAY_MIS_SKIP = 2048,
297         PATH_RAY_DIFFUSE_ANCESTOR = 4096,
298         PATH_RAY_SINGLE_PASS_DONE = 8192,
299
300         /* we need layer member flags to be the 20 upper bits */
301         PATH_RAY_LAYER_SHIFT = (32-20)
302 };
303
304 /* Closure Label */
305
306 typedef enum ClosureLabel {
307         LABEL_NONE = 0,
308         LABEL_TRANSMIT = 1,
309         LABEL_REFLECT = 2,
310         LABEL_DIFFUSE = 4,
311         LABEL_GLOSSY = 8,
312         LABEL_SINGULAR = 16,
313         LABEL_TRANSPARENT = 32,
314         LABEL_VOLUME_SCATTER = 64,
315 } ClosureLabel;
316
317 /* Render Passes */
318
319 typedef enum PassType {
320         PASS_NONE = 0,
321         PASS_COMBINED = (1 << 0),
322         PASS_DEPTH = (1 << 1),
323         PASS_NORMAL = (1 << 2),
324         PASS_UV = (1 << 3),
325         PASS_OBJECT_ID = (1 << 4),
326         PASS_MATERIAL_ID = (1 << 5),
327         PASS_DIFFUSE_COLOR = (1 << 6),
328         PASS_GLOSSY_COLOR = (1 << 7),
329         PASS_TRANSMISSION_COLOR = (1 << 8),
330         PASS_DIFFUSE_INDIRECT = (1 << 9),
331         PASS_GLOSSY_INDIRECT = (1 << 10),
332         PASS_TRANSMISSION_INDIRECT = (1 << 11),
333         PASS_DIFFUSE_DIRECT = (1 << 12),
334         PASS_GLOSSY_DIRECT = (1 << 13),
335         PASS_TRANSMISSION_DIRECT = (1 << 14),
336         PASS_EMISSION = (1 << 15),
337         PASS_BACKGROUND = (1 << 16),
338         PASS_AO = (1 << 17),
339         PASS_SHADOW = (1 << 18),
340         PASS_MOTION = (1 << 19),
341         PASS_MOTION_WEIGHT = (1 << 20),
342         PASS_MIST = (1 << 21),
343         PASS_SUBSURFACE_DIRECT = (1 << 22),
344         PASS_SUBSURFACE_INDIRECT = (1 << 23),
345         PASS_SUBSURFACE_COLOR = (1 << 24),
346         PASS_LIGHT = (1 << 25), /* no real pass, used to force use_light_pass */
347 #ifdef __KERNEL_DEBUG__
348         PASS_BVH_TRAVERSAL_STEPS = (1 << 26),
349 #endif
350 } PassType;
351
352 #define PASS_ALL (~0)
353
354 #ifdef __PASSES__
355
356 typedef ccl_addr_space struct PathRadiance {
357         int use_light_pass;
358
359         float3 emission;
360         float3 background;
361         float3 ao;
362
363         float3 indirect;
364         float3 direct_throughput;
365         float3 direct_emission;
366
367         float3 color_diffuse;
368         float3 color_glossy;
369         float3 color_transmission;
370         float3 color_subsurface;
371         float3 color_scatter;
372
373         float3 direct_diffuse;
374         float3 direct_glossy;
375         float3 direct_transmission;
376         float3 direct_subsurface;
377         float3 direct_scatter;
378
379         float3 indirect_diffuse;
380         float3 indirect_glossy;
381         float3 indirect_transmission;
382         float3 indirect_subsurface;
383         float3 indirect_scatter;
384
385         float3 path_diffuse;
386         float3 path_glossy;
387         float3 path_transmission;
388         float3 path_subsurface;
389         float3 path_scatter;
390
391         float4 shadow;
392         float mist;
393 } PathRadiance;
394
395 typedef struct BsdfEval {
396         int use_light_pass;
397
398         float3 diffuse;
399         float3 glossy;
400         float3 transmission;
401         float3 transparent;
402         float3 subsurface;
403         float3 scatter;
404 } BsdfEval;
405
406 #else
407
408 typedef ccl_addr_space float3 PathRadiance;
409 typedef float3 BsdfEval;
410
411 #endif
412
413 /* Shader Flag */
414
415 typedef enum ShaderFlag {
416         SHADER_SMOOTH_NORMAL = (1 << 31),
417         SHADER_CAST_SHADOW = (1 << 30),
418         SHADER_AREA_LIGHT = (1 << 29),
419         SHADER_USE_MIS = (1 << 28),
420         SHADER_EXCLUDE_DIFFUSE = (1 << 27),
421         SHADER_EXCLUDE_GLOSSY = (1 << 26),
422         SHADER_EXCLUDE_TRANSMIT = (1 << 25),
423         SHADER_EXCLUDE_CAMERA = (1 << 24),
424         SHADER_EXCLUDE_SCATTER = (1 << 23),
425         SHADER_EXCLUDE_ANY = (SHADER_EXCLUDE_DIFFUSE|SHADER_EXCLUDE_GLOSSY|SHADER_EXCLUDE_TRANSMIT|SHADER_EXCLUDE_CAMERA|SHADER_EXCLUDE_SCATTER),
426
427         SHADER_MASK = ~(SHADER_SMOOTH_NORMAL|SHADER_CAST_SHADOW|SHADER_AREA_LIGHT|SHADER_USE_MIS|SHADER_EXCLUDE_ANY)
428 } ShaderFlag;
429
430 /* Light Type */
431
432 typedef enum LightType {
433         LIGHT_POINT,
434         LIGHT_DISTANT,
435         LIGHT_BACKGROUND,
436         LIGHT_AREA,
437         LIGHT_SPOT,
438         LIGHT_TRIANGLE
439 } LightType;
440
441 /* Camera Type */
442
443 enum CameraType {
444         CAMERA_PERSPECTIVE,
445         CAMERA_ORTHOGRAPHIC,
446         CAMERA_PANORAMA
447 };
448
449 /* Panorama Type */
450
451 enum PanoramaType {
452         PANORAMA_EQUIRECTANGULAR,
453         PANORAMA_MIRRORBALL,
454         PANORAMA_FISHEYE_EQUIDISTANT,
455         PANORAMA_FISHEYE_EQUISOLID
456 };
457
458 /* Differential */
459
460 typedef struct differential3 {
461         float3 dx;
462         float3 dy;
463 } differential3;
464
465 typedef struct differential {
466         float dx;
467         float dy;
468 } differential;
469
470 /* Ray */
471
472 typedef struct Ray {
473 /* TODO(sergey): This is only needed because current AMD
474  * compilet has hard time bulding the kernel with this
475  * reshuffle. And at the same time reshuffle will cause
476  * less optimal CPU code in certain places.
477  *
478  * We'll get rid of this nasty eception once AMD compiler
479  * is fixed.
480  */
481 #ifndef __KERNEL_OPENCL_AMD__
482         float3 P;               /* origin */
483         float3 D;               /* direction */
484
485         float t;                /* length of the ray */
486         float time;             /* time (for motion blur) */
487 #else
488         float t;                /* length of the ray */
489         float time;             /* time (for motion blur) */
490         float3 P;               /* origin */
491         float3 D;               /* direction */
492 #endif
493
494 #ifdef __RAY_DIFFERENTIALS__
495         differential3 dP;
496         differential3 dD;
497 #endif
498 } Ray;
499
500 /* Intersection */
501
502 typedef ccl_addr_space struct Intersection {
503         float t, u, v;
504         int prim;
505         int object;
506         int type;
507
508 #ifdef __KERNEL_DEBUG__
509         int num_traversal_steps;
510 #endif
511 } Intersection;
512
513 /* Primitives */
514
515 typedef enum PrimitiveType {
516         PRIMITIVE_NONE = 0,
517         PRIMITIVE_TRIANGLE = 1,
518         PRIMITIVE_MOTION_TRIANGLE = 2,
519         PRIMITIVE_CURVE = 4,
520         PRIMITIVE_MOTION_CURVE = 8,
521
522         PRIMITIVE_ALL_TRIANGLE = (PRIMITIVE_TRIANGLE|PRIMITIVE_MOTION_TRIANGLE),
523         PRIMITIVE_ALL_CURVE = (PRIMITIVE_CURVE|PRIMITIVE_MOTION_CURVE),
524         PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE|PRIMITIVE_MOTION_CURVE),
525         PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE|PRIMITIVE_ALL_CURVE),
526
527         /* Total number of different primitives.
528          * NOTE: This is an actual value, not a bitflag.
529          */
530         PRIMITIVE_NUM_TOTAL = 4,
531 } PrimitiveType;
532
533 #define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << 16) | type)
534 #define PRIMITIVE_UNPACK_SEGMENT(type) (type >> 16)
535
536 /* Attributes */
537
538 #define ATTR_PRIM_TYPES         2
539 #define ATTR_PRIM_CURVE         1
540
541 typedef enum AttributeElement {
542         ATTR_ELEMENT_NONE,
543         ATTR_ELEMENT_OBJECT,
544         ATTR_ELEMENT_MESH,
545         ATTR_ELEMENT_FACE,
546         ATTR_ELEMENT_VERTEX,
547         ATTR_ELEMENT_VERTEX_MOTION,
548         ATTR_ELEMENT_CORNER,
549         ATTR_ELEMENT_CORNER_BYTE,
550         ATTR_ELEMENT_CURVE,
551         ATTR_ELEMENT_CURVE_KEY,
552         ATTR_ELEMENT_CURVE_KEY_MOTION,
553         ATTR_ELEMENT_VOXEL
554 } AttributeElement;
555
556 typedef enum AttributeStandard {
557         ATTR_STD_NONE = 0,
558         ATTR_STD_VERTEX_NORMAL,
559         ATTR_STD_FACE_NORMAL,
560         ATTR_STD_UV,
561         ATTR_STD_UV_TANGENT,
562         ATTR_STD_UV_TANGENT_SIGN,
563         ATTR_STD_GENERATED,
564         ATTR_STD_GENERATED_TRANSFORM,
565         ATTR_STD_POSITION_UNDEFORMED,
566         ATTR_STD_POSITION_UNDISPLACED,
567         ATTR_STD_MOTION_VERTEX_POSITION,
568         ATTR_STD_MOTION_VERTEX_NORMAL,
569         ATTR_STD_PARTICLE,
570         ATTR_STD_CURVE_INTERCEPT,
571         ATTR_STD_PTEX_FACE_ID,
572         ATTR_STD_PTEX_UV,
573         ATTR_STD_VOLUME_DENSITY,
574         ATTR_STD_VOLUME_COLOR,
575         ATTR_STD_VOLUME_FLAME,
576         ATTR_STD_VOLUME_HEAT,
577         ATTR_STD_VOLUME_VELOCITY,
578         ATTR_STD_POINTINESS,
579         ATTR_STD_NUM,
580
581         ATTR_STD_NOT_FOUND = ~0
582 } AttributeStandard;
583
584 /* Closure data */
585
586 #ifdef __MULTI_CLOSURE__
587 #  ifndef __MAX_CLOSURE__
588 #     define MAX_CLOSURE 64
589 #  else
590 #    define MAX_CLOSURE __MAX_CLOSURE__
591 #  endif
592 #else
593 #define MAX_CLOSURE 1
594 #endif
595
596 /* This struct is to be 16 bytes aligned, we also keep some extra precautions:
597  * - All the float3 members are in the beginning of the struct, so compiler
598  *   does not put own padding trying to align this members.
599  * - We make sure OSL pointer is also 16 bytes aligned.
600  */
601 typedef ccl_addr_space struct ShaderClosure {
602         float3 weight;
603         float3 N;
604         float3 T;
605
606         ClosureType type;
607         float sample_weight;
608         float data0;
609         float data1;
610         float data2;
611         int pad1, pad2, pad3;
612
613 #ifdef __OSL__
614         void *prim, *pad4;
615 #endif
616 } ShaderClosure;
617
618 /* Shader Context
619  *
620  * For OSL we recycle a fixed number of contexts for speed */
621
622 typedef enum ShaderContext {
623         SHADER_CONTEXT_MAIN = 0,
624         SHADER_CONTEXT_INDIRECT = 1,
625         SHADER_CONTEXT_EMISSION = 2,
626         SHADER_CONTEXT_SHADOW = 3,
627         SHADER_CONTEXT_SSS = 4,
628         SHADER_CONTEXT_VOLUME = 5,
629         SHADER_CONTEXT_NUM = 6
630 } ShaderContext;
631
632 /* Shader Data
633  *
634  * Main shader state at a point on the surface or in a volume. All coordinates
635  * are in world space. */
636
637 enum ShaderDataFlag {
638         /* runtime flags */
639         SD_BACKFACING     = (1 << 0),   /* backside of surface? */
640         SD_EMISSION       = (1 << 1),   /* have emissive closure? */
641         SD_BSDF           = (1 << 2),   /* have bsdf closure? */
642         SD_BSDF_HAS_EVAL  = (1 << 3),   /* have non-singular bsdf closure? */
643         SD_BSSRDF         = (1 << 4),   /* have bssrdf */
644         SD_HOLDOUT        = (1 << 5),   /* have holdout closure? */
645         SD_ABSORPTION     = (1 << 6),   /* have volume absorption closure? */
646         SD_SCATTER        = (1 << 7),   /* have volume phase closure? */
647         SD_AO             = (1 << 8),   /* have ao closure? */
648         SD_TRANSPARENT    = (1 << 9),  /* have transparent closure? */
649
650         SD_CLOSURE_FLAGS = (SD_EMISSION|SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSSRDF|
651                             SD_HOLDOUT|SD_ABSORPTION|SD_SCATTER|SD_AO),
652
653         /* shader flags */
654         SD_USE_MIS                = (1 << 10),  /* direct light sample */
655         SD_HAS_TRANSPARENT_SHADOW = (1 << 11),  /* has transparent shadow */
656         SD_HAS_VOLUME             = (1 << 12),  /* has volume shader */
657         SD_HAS_ONLY_VOLUME        = (1 << 13),  /* has only volume shader, no surface */
658         SD_HETEROGENEOUS_VOLUME   = (1 << 14),  /* has heterogeneous volume */
659         SD_HAS_BSSRDF_BUMP        = (1 << 15),  /* bssrdf normal uses bump */
660         SD_VOLUME_EQUIANGULAR     = (1 << 16),  /* use equiangular sampling */
661         SD_VOLUME_MIS             = (1 << 17),  /* use multiple importance sampling */
662         SD_VOLUME_CUBIC           = (1 << 18),  /* use cubic interpolation for voxels */
663         SD_HAS_BUMP               = (1 << 19),  /* has data connected to the displacement input */
664
665         SD_SHADER_FLAGS = (SD_USE_MIS|SD_HAS_TRANSPARENT_SHADOW|SD_HAS_VOLUME|
666                            SD_HAS_ONLY_VOLUME|SD_HETEROGENEOUS_VOLUME|
667                            SD_HAS_BSSRDF_BUMP|SD_VOLUME_EQUIANGULAR|SD_VOLUME_MIS|
668                            SD_VOLUME_CUBIC|SD_HAS_BUMP),
669
670         /* object flags */
671         SD_HOLDOUT_MASK             = (1 << 20),  /* holdout for camera rays */
672         SD_OBJECT_MOTION            = (1 << 21),  /* has object motion blur */
673         SD_TRANSFORM_APPLIED        = (1 << 22),  /* vertices have transform applied */
674         SD_NEGATIVE_SCALE_APPLIED   = (1 << 23),  /* vertices have negative scale applied */
675         SD_OBJECT_HAS_VOLUME        = (1 << 24),  /* object has a volume shader */
676         SD_OBJECT_INTERSECTS_VOLUME = (1 << 25),  /* object intersects AABB of an object with volume shader */
677         SD_OBJECT_HAS_VERTEX_MOTION = (1 << 26),  /* has position for motion vertices */
678
679         SD_OBJECT_FLAGS = (SD_HOLDOUT_MASK|SD_OBJECT_MOTION|SD_TRANSFORM_APPLIED|
680                            SD_NEGATIVE_SCALE_APPLIED|SD_OBJECT_HAS_VOLUME|
681                            SD_OBJECT_INTERSECTS_VOLUME)
682 };
683
684 struct KernelGlobals;
685
686 #ifdef __SPLIT_KERNEL__
687 #define SD_VAR(type, what) ccl_global type *what;
688 #define SD_CLOSURE_VAR(type, what, max_closure) type *what;
689 #define TIDX (get_global_id(1) * get_global_size(0) + get_global_id(0))
690 #define ccl_fetch(s, t) (s->t[TIDX])
691 #define ccl_fetch_array(s, t, index) (&s->t[TIDX * MAX_CLOSURE + index])
692 #else
693 #define SD_VAR(type, what) type what;
694 #define SD_CLOSURE_VAR(type, what, max_closure) type what[max_closure];
695 #define ccl_fetch(s, t) (s->t)
696 #define ccl_fetch_array(s, t, index) (&s->t[index])
697 #endif
698
699 typedef ccl_addr_space struct ShaderData {
700
701 #include "kernel_shaderdata_vars.h"
702
703 } ShaderData;
704
705 /* Path State */
706
707 #ifdef __VOLUME__
708 typedef struct VolumeStack {
709         int object;
710         int shader;
711 } VolumeStack;
712 #endif
713
714 typedef struct PathState {
715         /* see enum PathRayFlag */
716         int flag;          
717
718         /* random number generator state */
719         int rng_offset;                 /* dimension offset */
720         int rng_offset_bsdf;    /* dimension offset for picking bsdf */
721         int sample;                     /* path sample number */
722         int num_samples;                /* total number of times this path will be sampled */
723
724         /* bounce counting */
725         int bounce;
726         int diffuse_bounce;
727         int glossy_bounce;
728         int transmission_bounce;
729         int transparent_bounce;
730
731         /* multiple importance sampling */
732         float min_ray_pdf; /* smallest bounce pdf over entire path up to now */
733         float ray_pdf;     /* last bounce pdf */
734 #ifdef __LAMP_MIS__
735         float ray_t;       /* accumulated distance through transparent surfaces */
736 #endif
737
738         /* volume rendering */
739 #ifdef __VOLUME__
740         int volume_bounce;
741         RNG rng_congruential;
742         VolumeStack volume_stack[VOLUME_STACK_SIZE];
743 #endif
744 } PathState;
745
746 /* Constant Kernel Data
747  *
748  * These structs are passed from CPU to various devices, and the struct layout
749  * must match exactly. Structs are padded to ensure 16 byte alignment, and we
750  * do not use float3 because its size may not be the same on all devices. */
751
752 typedef struct KernelCamera {
753         /* type */
754         int type;
755
756         /* panorama */
757         int panorama_type;
758         float fisheye_fov;
759         float fisheye_lens;
760         float4 equirectangular_range;
761
762         /* matrices */
763         Transform cameratoworld;
764         Transform rastertocamera;
765
766         /* differentials */
767         float4 dx;
768         float4 dy;
769
770         /* depth of field */
771         float aperturesize;
772         float blades;
773         float bladesrotation;
774         float focaldistance;
775
776         /* motion blur */
777         float shuttertime;
778         int have_motion;
779
780         /* clipping */
781         float nearclip;
782         float cliplength;
783
784         /* sensor size */
785         float sensorwidth;
786         float sensorheight;
787
788         /* render size */
789         float width, height;
790         int resolution;
791
792         /* anamorphic lens bokeh */
793         float inv_aperture_ratio;
794
795         int is_inside_volume;
796         int pad2;
797
798         /* more matrices */
799         Transform screentoworld;
800         Transform rastertoworld;
801         /* work around cuda sm 2.0 crash, this seems to
802          * cross some limit in combination with motion 
803          * Transform ndctoworld; */
804         Transform worldtoscreen;
805         Transform worldtoraster;
806         Transform worldtondc;
807         Transform worldtocamera;
808
809         MotionTransform motion;
810 } KernelCamera;
811
812 typedef struct KernelFilm {
813         float exposure;
814         int pass_flag;
815         int pass_stride;
816         int use_light_pass;
817
818         int pass_combined;
819         int pass_depth;
820         int pass_normal;
821         int pass_motion;
822
823         int pass_motion_weight;
824         int pass_uv;
825         int pass_object_id;
826         int pass_material_id;
827
828         int pass_diffuse_color;
829         int pass_glossy_color;
830         int pass_transmission_color;
831         int pass_subsurface_color;
832         
833         int pass_diffuse_indirect;
834         int pass_glossy_indirect;
835         int pass_transmission_indirect;
836         int pass_subsurface_indirect;
837         
838         int pass_diffuse_direct;
839         int pass_glossy_direct;
840         int pass_transmission_direct;
841         int pass_subsurface_direct;
842         
843         int pass_emission;
844         int pass_background;
845         int pass_ao;
846         float pass_alpha_threshold;
847
848         int pass_shadow;
849         float pass_shadow_scale;
850         int filter_table_offset;
851         int pass_pad2;
852
853         int pass_mist;
854         float mist_start;
855         float mist_inv_depth;
856         float mist_falloff;
857
858 #ifdef __KERNEL_DEBUG__
859         int pass_bvh_traversal_steps;
860         int pass_pad3, pass_pad4, pass_pad5;
861 #endif
862 } KernelFilm;
863
864 typedef struct KernelBackground {
865         /* only shader index */
866         int surface_shader;
867         int volume_shader;
868         int transparent;
869         int pad;
870
871         /* ambient occlusion */
872         float ao_factor;
873         float ao_distance;
874         float ao_pad1, ao_pad2;
875 } KernelBackground;
876
877 typedef struct KernelIntegrator {
878         /* emission */
879         int use_direct_light;
880         int use_ambient_occlusion;
881         int num_distribution;
882         int num_all_lights;
883         float pdf_triangles;
884         float pdf_lights;
885         float inv_pdf_lights;
886         int pdf_background_res;
887
888         /* light portals */
889         float portal_pdf;
890         int num_portals;
891         int portal_offset;
892
893         /* bounces */
894         int min_bounce;
895         int max_bounce;
896
897         int max_diffuse_bounce;
898         int max_glossy_bounce;
899         int max_transmission_bounce;
900         int max_volume_bounce;
901
902         /* transparent */
903         int transparent_min_bounce;
904         int transparent_max_bounce;
905         int transparent_shadows;
906
907         /* caustics */
908         int caustics_reflective;
909         int caustics_refractive;
910         float filter_glossy;
911
912         /* seed */
913         int seed;
914
915         /* render layer */
916         int layer_flag;
917
918         /* clamp */
919         float sample_clamp_direct;
920         float sample_clamp_indirect;
921
922         /* branched path */
923         int branched;
924         int diffuse_samples;
925         int glossy_samples;
926         int transmission_samples;
927         int ao_samples;
928         int mesh_light_samples;
929         int subsurface_samples;
930         int sample_all_lights_direct;
931         int sample_all_lights_indirect;
932
933         /* mis */
934         int use_lamp_mis;
935
936         /* sampler */
937         int sampling_pattern;
938         int aa_samples;
939
940         /* volume render */
941         int use_volumes;
942         int volume_max_steps;
943         float volume_step_size;
944         int volume_samples;
945
946         int pad;
947 } KernelIntegrator;
948
949 typedef struct KernelBVH {
950         /* root node */
951         int root;
952         int attributes_map_stride;
953         int have_motion;
954         int have_curves;
955         int have_instancing;
956         int use_qbvh;
957         int pad1, pad2;
958 } KernelBVH;
959
960 typedef enum CurveFlag {
961         /* runtime flags */
962         CURVE_KN_BACKFACING = 1,                                /* backside of cylinder? */
963         CURVE_KN_ENCLOSEFILTER = 2,                             /* don't consider strands surrounding start point? */
964         CURVE_KN_INTERPOLATE = 4,                               /* render as a curve? */
965         CURVE_KN_ACCURATE = 8,                                  /* use accurate intersections test? */
966         CURVE_KN_INTERSECTCORRECTION = 16,              /* correct for width after determing closest midpoint? */
967         CURVE_KN_TRUETANGENTGNORMAL = 32,               /* use tangent normal for geometry? */
968         CURVE_KN_RIBBONS = 64,                                  /* use flat curve ribbons */
969 } CurveFlag;
970
971 typedef struct KernelCurves {
972         int curveflags;
973         int subdivisions;
974
975         float minimum_width;
976         float maximum_width;
977 } KernelCurves;
978
979 typedef struct KernelTables {
980         int beckmann_offset;
981         int pad1, pad2, pad3;
982 } KernelTables;
983
984 typedef struct KernelData {
985         KernelCamera cam;
986         KernelFilm film;
987         KernelBackground background;
988         KernelIntegrator integrator;
989         KernelBVH bvh;
990         KernelCurves curve;
991         KernelTables tables;
992 } KernelData;
993
994 #ifdef __KERNEL_DEBUG__
995 typedef ccl_addr_space struct DebugData {
996         // Total number of BVH node traversal steps and primitives intersections
997         // for the camera rays.
998         int num_bvh_traversal_steps;
999 } DebugData;
1000 #endif
1001
1002 /* Declarations required for split kernel */
1003
1004 /* Macro for queues */
1005 /* Value marking queue's empty slot */
1006 #define QUEUE_EMPTY_SLOT -1
1007
1008 /*
1009 * Queue 1 - Active rays
1010 * Queue 2 - Background queue
1011 * Queue 3 - Shadow ray cast kernel - AO
1012 * Queeu 4 - Shadow ray cast kernel - direct lighting
1013 */
1014 #define NUM_QUEUES 4
1015
1016 /* Queue names */
1017 enum QueueNumber {
1018         QUEUE_ACTIVE_AND_REGENERATED_RAYS,         /* All active rays and regenerated rays are enqueued here */
1019         QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS,      /* All
1020                                                     * 1.Background-hit rays,
1021                                                     * 2.Rays that has exited path-iteration but needs to update output buffer
1022                                                     * 3.Rays to be regenerated
1023                                                     * are enqueued here */
1024         QUEUE_SHADOW_RAY_CAST_AO_RAYS,             /* All rays for which a shadow ray should be cast to determine radiance
1025                                                       contribution for AO are enqueued here */
1026         QUEUE_SHADOW_RAY_CAST_DL_RAYS,             /* All rays for which a shadow ray should be cast to determine radiance
1027                                                       contributuin for direct lighting are enqueued here */
1028 };
1029
1030 /* We use RAY_STATE_MASK to get ray_state (enums 0 to 5) */
1031 #define RAY_STATE_MASK 0x007
1032 #define RAY_FLAG_MASK 0x0F8
1033 enum RayState {
1034         RAY_ACTIVE = 0,             // Denotes ray is actively involved in path-iteration
1035         RAY_INACTIVE = 1,           // Denotes ray has completed processing all samples and is inactive
1036         RAY_UPDATE_BUFFER = 2,      // Denoted ray has exited path-iteration and needs to update output buffer
1037         RAY_HIT_BACKGROUND = 3,     // Donotes ray has hit background
1038         RAY_TO_REGENERATE = 4,      // Denotes ray has to be regenerated
1039         RAY_REGENERATED = 5,        // Denotes ray has been regenerated
1040         RAY_SKIP_DL = 6,            // Denotes ray should skip direct lighting
1041         RAY_SHADOW_RAY_CAST_AO = 16, // Flag's ray has to execute shadow blocked function in AO part
1042         RAY_SHADOW_RAY_CAST_DL = 32 // Flag's ray has to execute shadow blocked function in direct lighting part
1043 };
1044
1045 #define ASSIGN_RAY_STATE(ray_state, ray_index, state) (ray_state[ray_index] = ((ray_state[ray_index] & RAY_FLAG_MASK) | state))
1046 #define IS_STATE(ray_state, ray_index, state) ((ray_state[ray_index] & RAY_STATE_MASK) == state)
1047 #define ADD_RAY_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] = (ray_state[ray_index] | flag))
1048 #define REMOVE_RAY_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] = (ray_state[ray_index] & (~flag)))
1049 #define IS_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] & flag)
1050
1051 CCL_NAMESPACE_END
1052
1053 #endif /*  __KERNEL_TYPES_H__ */
1054