Cycles: OpenCL kernel split
[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         float3 P;               /* origin */
474         float3 D;               /* direction */
475
476         float t;                /* length of the ray */
477         float time;             /* time (for motion blur) */
478 #ifdef __RAY_DIFFERENTIALS__
479         differential3 dP;
480         differential3 dD;
481 #endif
482 } Ray;
483
484 /* Intersection */
485
486 typedef ccl_addr_space struct Intersection {
487         float t, u, v;
488         int prim;
489         int object;
490         int type;
491
492 #ifdef __KERNEL_DEBUG__
493         int num_traversal_steps;
494 #endif
495 } Intersection;
496
497 /* Primitives */
498
499 typedef enum PrimitiveType {
500         PRIMITIVE_NONE = 0,
501         PRIMITIVE_TRIANGLE = 1,
502         PRIMITIVE_MOTION_TRIANGLE = 2,
503         PRIMITIVE_CURVE = 4,
504         PRIMITIVE_MOTION_CURVE = 8,
505
506         PRIMITIVE_ALL_TRIANGLE = (PRIMITIVE_TRIANGLE|PRIMITIVE_MOTION_TRIANGLE),
507         PRIMITIVE_ALL_CURVE = (PRIMITIVE_CURVE|PRIMITIVE_MOTION_CURVE),
508         PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE|PRIMITIVE_MOTION_CURVE),
509         PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE|PRIMITIVE_ALL_CURVE),
510
511         /* Total number of different primitives.
512          * NOTE: This is an actual value, not a bitflag.
513          */
514         PRIMITIVE_NUM_TOTAL = 4,
515 } PrimitiveType;
516
517 #define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << 16) | type)
518 #define PRIMITIVE_UNPACK_SEGMENT(type) (type >> 16)
519
520 /* Attributes */
521
522 #define ATTR_PRIM_TYPES         2
523 #define ATTR_PRIM_CURVE         1
524
525 typedef enum AttributeElement {
526         ATTR_ELEMENT_NONE,
527         ATTR_ELEMENT_OBJECT,
528         ATTR_ELEMENT_MESH,
529         ATTR_ELEMENT_FACE,
530         ATTR_ELEMENT_VERTEX,
531         ATTR_ELEMENT_VERTEX_MOTION,
532         ATTR_ELEMENT_CORNER,
533         ATTR_ELEMENT_CORNER_BYTE,
534         ATTR_ELEMENT_CURVE,
535         ATTR_ELEMENT_CURVE_KEY,
536         ATTR_ELEMENT_CURVE_KEY_MOTION,
537         ATTR_ELEMENT_VOXEL
538 } AttributeElement;
539
540 typedef enum AttributeStandard {
541         ATTR_STD_NONE = 0,
542         ATTR_STD_VERTEX_NORMAL,
543         ATTR_STD_FACE_NORMAL,
544         ATTR_STD_UV,
545         ATTR_STD_UV_TANGENT,
546         ATTR_STD_UV_TANGENT_SIGN,
547         ATTR_STD_GENERATED,
548         ATTR_STD_GENERATED_TRANSFORM,
549         ATTR_STD_POSITION_UNDEFORMED,
550         ATTR_STD_POSITION_UNDISPLACED,
551         ATTR_STD_MOTION_VERTEX_POSITION,
552         ATTR_STD_MOTION_VERTEX_NORMAL,
553         ATTR_STD_PARTICLE,
554         ATTR_STD_CURVE_INTERCEPT,
555         ATTR_STD_PTEX_FACE_ID,
556         ATTR_STD_PTEX_UV,
557         ATTR_STD_VOLUME_DENSITY,
558         ATTR_STD_VOLUME_COLOR,
559         ATTR_STD_VOLUME_FLAME,
560         ATTR_STD_VOLUME_HEAT,
561         ATTR_STD_VOLUME_VELOCITY,
562         ATTR_STD_POINTINESS,
563         ATTR_STD_NUM,
564
565         ATTR_STD_NOT_FOUND = ~0
566 } AttributeStandard;
567
568 /* Closure data */
569
570 #ifdef __MULTI_CLOSURE__
571 #  ifndef __MAX_CLOSURE__
572 #     define MAX_CLOSURE 64
573 #  else
574 #    define MAX_CLOSURE __MAX_CLOSURE__
575 #  endif
576 #else
577 #define MAX_CLOSURE 1
578 #endif
579
580 /* This struct is to be 16 bytes aligned, we also keep some extra precautions:
581  * - All the float3 members are in the beginning of the struct, so compiler
582  *   does not put own padding trying to align this members.
583  * - We make sure OSL pointer is also 16 bytes aligned.
584  */
585 typedef ccl_addr_space struct ShaderClosure {
586         float3 weight;
587         float3 N;
588         float3 T;
589
590         ClosureType type;
591         float sample_weight;
592         float data0;
593         float data1;
594         float data2;
595         int pad1, pad2, pad3;
596
597 #ifdef __OSL__
598         void *prim, *pad4;
599 #endif
600 } ShaderClosure;
601
602 /* Shader Context
603  *
604  * For OSL we recycle a fixed number of contexts for speed */
605
606 typedef enum ShaderContext {
607         SHADER_CONTEXT_MAIN = 0,
608         SHADER_CONTEXT_INDIRECT = 1,
609         SHADER_CONTEXT_EMISSION = 2,
610         SHADER_CONTEXT_SHADOW = 3,
611         SHADER_CONTEXT_SSS = 4,
612         SHADER_CONTEXT_VOLUME = 5,
613         SHADER_CONTEXT_NUM = 6
614 } ShaderContext;
615
616 /* Shader Data
617  *
618  * Main shader state at a point on the surface or in a volume. All coordinates
619  * are in world space. */
620
621 enum ShaderDataFlag {
622         /* runtime flags */
623         SD_BACKFACING     = (1 << 0),   /* backside of surface? */
624         SD_EMISSION       = (1 << 1),   /* have emissive closure? */
625         SD_BSDF           = (1 << 2),   /* have bsdf closure? */
626         SD_BSDF_HAS_EVAL  = (1 << 3),   /* have non-singular bsdf closure? */
627         SD_BSSRDF         = (1 << 4),   /* have bssrdf */
628         SD_HOLDOUT        = (1 << 5),   /* have holdout closure? */
629         SD_ABSORPTION     = (1 << 6),   /* have volume absorption closure? */
630         SD_SCATTER        = (1 << 7),   /* have volume phase closure? */
631         SD_AO             = (1 << 8),   /* have ao closure? */
632         SD_TRANSPARENT    = (1 << 9),  /* have transparent closure? */
633
634         SD_CLOSURE_FLAGS = (SD_EMISSION|SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSSRDF|
635                             SD_HOLDOUT|SD_ABSORPTION|SD_SCATTER|SD_AO),
636
637         /* shader flags */
638         SD_USE_MIS                = (1 << 10),  /* direct light sample */
639         SD_HAS_TRANSPARENT_SHADOW = (1 << 11),  /* has transparent shadow */
640         SD_HAS_VOLUME             = (1 << 12),  /* has volume shader */
641         SD_HAS_ONLY_VOLUME        = (1 << 13),  /* has only volume shader, no surface */
642         SD_HETEROGENEOUS_VOLUME   = (1 << 14),  /* has heterogeneous volume */
643         SD_HAS_BSSRDF_BUMP        = (1 << 15),  /* bssrdf normal uses bump */
644         SD_VOLUME_EQUIANGULAR     = (1 << 16),  /* use equiangular sampling */
645         SD_VOLUME_MIS             = (1 << 17),  /* use multiple importance sampling */
646         SD_VOLUME_CUBIC           = (1 << 18),  /* use cubic interpolation for voxels */
647         SD_HAS_BUMP               = (1 << 19),  /* has data connected to the displacement input */
648
649         SD_SHADER_FLAGS = (SD_USE_MIS|SD_HAS_TRANSPARENT_SHADOW|SD_HAS_VOLUME|
650                            SD_HAS_ONLY_VOLUME|SD_HETEROGENEOUS_VOLUME|
651                            SD_HAS_BSSRDF_BUMP|SD_VOLUME_EQUIANGULAR|SD_VOLUME_MIS|
652                            SD_VOLUME_CUBIC|SD_HAS_BUMP),
653
654         /* object flags */
655         SD_HOLDOUT_MASK             = (1 << 20),  /* holdout for camera rays */
656         SD_OBJECT_MOTION            = (1 << 21),  /* has object motion blur */
657         SD_TRANSFORM_APPLIED        = (1 << 22),  /* vertices have transform applied */
658         SD_NEGATIVE_SCALE_APPLIED   = (1 << 23),  /* vertices have negative scale applied */
659         SD_OBJECT_HAS_VOLUME        = (1 << 24),  /* object has a volume shader */
660         SD_OBJECT_INTERSECTS_VOLUME = (1 << 25),  /* object intersects AABB of an object with volume shader */
661         SD_OBJECT_HAS_VERTEX_MOTION = (1 << 26),  /* has position for motion vertices */
662
663         SD_OBJECT_FLAGS = (SD_HOLDOUT_MASK|SD_OBJECT_MOTION|SD_TRANSFORM_APPLIED|
664                            SD_NEGATIVE_SCALE_APPLIED|SD_OBJECT_HAS_VOLUME|
665                            SD_OBJECT_INTERSECTS_VOLUME)
666 };
667
668 struct KernelGlobals;
669
670 #ifdef __SPLIT_KERNEL__
671 #define SD_VAR(type, what) ccl_global type *what;
672 #define SD_CLOSURE_VAR(type, what, max_closure) type *what;
673 #define TIDX (get_global_id(1) * get_global_size(0) + get_global_id(0))
674 #define ccl_fetch(s, t) (s->t[TIDX])
675 #define ccl_fetch_array(s, t, index) (&s->t[TIDX * MAX_CLOSURE + index])
676 #else
677 #define SD_VAR(type, what) type what;
678 #define SD_CLOSURE_VAR(type, what, max_closure) type what[max_closure];
679 #define ccl_fetch(s, t) (s->t)
680 #define ccl_fetch_array(s, t, index) (&s->t[index])
681 #endif
682
683 typedef ccl_addr_space struct ShaderData {
684
685 #include "kernel_shaderdata_vars.h"
686
687 } ShaderData;
688
689 /* Path State */
690
691 #ifdef __VOLUME__
692 typedef struct VolumeStack {
693         int object;
694         int shader;
695 } VolumeStack;
696 #endif
697
698 typedef struct PathState {
699         /* see enum PathRayFlag */
700         int flag;          
701
702         /* random number generator state */
703         int rng_offset;                 /* dimension offset */
704         int rng_offset_bsdf;    /* dimension offset for picking bsdf */
705         int sample;                     /* path sample number */
706         int num_samples;                /* total number of times this path will be sampled */
707
708         /* bounce counting */
709         int bounce;
710         int diffuse_bounce;
711         int glossy_bounce;
712         int transmission_bounce;
713         int transparent_bounce;
714
715         /* multiple importance sampling */
716         float min_ray_pdf; /* smallest bounce pdf over entire path up to now */
717         float ray_pdf;     /* last bounce pdf */
718 #ifdef __LAMP_MIS__
719         float ray_t;       /* accumulated distance through transparent surfaces */
720 #endif
721
722         /* volume rendering */
723 #ifdef __VOLUME__
724         int volume_bounce;
725         RNG rng_congruential;
726         VolumeStack volume_stack[VOLUME_STACK_SIZE];
727 #endif
728 } PathState;
729
730 /* Constant Kernel Data
731  *
732  * These structs are passed from CPU to various devices, and the struct layout
733  * must match exactly. Structs are padded to ensure 16 byte alignment, and we
734  * do not use float3 because its size may not be the same on all devices. */
735
736 typedef struct KernelCamera {
737         /* type */
738         int type;
739
740         /* panorama */
741         int panorama_type;
742         float fisheye_fov;
743         float fisheye_lens;
744         float4 equirectangular_range;
745
746         /* matrices */
747         Transform cameratoworld;
748         Transform rastertocamera;
749
750         /* differentials */
751         float4 dx;
752         float4 dy;
753
754         /* depth of field */
755         float aperturesize;
756         float blades;
757         float bladesrotation;
758         float focaldistance;
759
760         /* motion blur */
761         float shuttertime;
762         int have_motion;
763
764         /* clipping */
765         float nearclip;
766         float cliplength;
767
768         /* sensor size */
769         float sensorwidth;
770         float sensorheight;
771
772         /* render size */
773         float width, height;
774         int resolution;
775
776         /* anamorphic lens bokeh */
777         float inv_aperture_ratio;
778
779         int is_inside_volume;
780         int pad2;
781
782         /* more matrices */
783         Transform screentoworld;
784         Transform rastertoworld;
785         /* work around cuda sm 2.0 crash, this seems to
786          * cross some limit in combination with motion 
787          * Transform ndctoworld; */
788         Transform worldtoscreen;
789         Transform worldtoraster;
790         Transform worldtondc;
791         Transform worldtocamera;
792
793         MotionTransform motion;
794 } KernelCamera;
795
796 typedef struct KernelFilm {
797         float exposure;
798         int pass_flag;
799         int pass_stride;
800         int use_light_pass;
801
802         int pass_combined;
803         int pass_depth;
804         int pass_normal;
805         int pass_motion;
806
807         int pass_motion_weight;
808         int pass_uv;
809         int pass_object_id;
810         int pass_material_id;
811
812         int pass_diffuse_color;
813         int pass_glossy_color;
814         int pass_transmission_color;
815         int pass_subsurface_color;
816         
817         int pass_diffuse_indirect;
818         int pass_glossy_indirect;
819         int pass_transmission_indirect;
820         int pass_subsurface_indirect;
821         
822         int pass_diffuse_direct;
823         int pass_glossy_direct;
824         int pass_transmission_direct;
825         int pass_subsurface_direct;
826         
827         int pass_emission;
828         int pass_background;
829         int pass_ao;
830         float pass_alpha_threshold;
831
832         int pass_shadow;
833         float pass_shadow_scale;
834         int filter_table_offset;
835         int pass_pad2;
836
837         int pass_mist;
838         float mist_start;
839         float mist_inv_depth;
840         float mist_falloff;
841
842 #ifdef __KERNEL_DEBUG__
843         int pass_bvh_traversal_steps;
844         int pass_pad3, pass_pad4, pass_pad5;
845 #endif
846 } KernelFilm;
847
848 typedef struct KernelBackground {
849         /* only shader index */
850         int surface_shader;
851         int volume_shader;
852         int transparent;
853         int pad;
854
855         /* ambient occlusion */
856         float ao_factor;
857         float ao_distance;
858         float ao_pad1, ao_pad2;
859 } KernelBackground;
860
861 typedef struct KernelIntegrator {
862         /* emission */
863         int use_direct_light;
864         int use_ambient_occlusion;
865         int num_distribution;
866         int num_all_lights;
867         float pdf_triangles;
868         float pdf_lights;
869         float inv_pdf_lights;
870         int pdf_background_res;
871
872         /* light portals */
873         float portal_pdf;
874         int num_portals;
875         int portal_offset;
876
877         /* bounces */
878         int min_bounce;
879         int max_bounce;
880
881         int max_diffuse_bounce;
882         int max_glossy_bounce;
883         int max_transmission_bounce;
884         int max_volume_bounce;
885
886         /* transparent */
887         int transparent_min_bounce;
888         int transparent_max_bounce;
889         int transparent_shadows;
890
891         /* caustics */
892         int caustics_reflective;
893         int caustics_refractive;
894         float filter_glossy;
895
896         /* seed */
897         int seed;
898
899         /* render layer */
900         int layer_flag;
901
902         /* clamp */
903         float sample_clamp_direct;
904         float sample_clamp_indirect;
905
906         /* branched path */
907         int branched;
908         int diffuse_samples;
909         int glossy_samples;
910         int transmission_samples;
911         int ao_samples;
912         int mesh_light_samples;
913         int subsurface_samples;
914         int sample_all_lights_direct;
915         int sample_all_lights_indirect;
916
917         /* mis */
918         int use_lamp_mis;
919
920         /* sampler */
921         int sampling_pattern;
922         int aa_samples;
923
924         /* volume render */
925         int use_volumes;
926         int volume_max_steps;
927         float volume_step_size;
928         int volume_samples;
929
930         int pad;
931 } KernelIntegrator;
932
933 typedef struct KernelBVH {
934         /* root node */
935         int root;
936         int attributes_map_stride;
937         int have_motion;
938         int have_curves;
939         int have_instancing;
940         int use_qbvh;
941         int pad1, pad2;
942 } KernelBVH;
943
944 typedef enum CurveFlag {
945         /* runtime flags */
946         CURVE_KN_BACKFACING = 1,                                /* backside of cylinder? */
947         CURVE_KN_ENCLOSEFILTER = 2,                             /* don't consider strands surrounding start point? */
948         CURVE_KN_INTERPOLATE = 4,                               /* render as a curve? */
949         CURVE_KN_ACCURATE = 8,                                  /* use accurate intersections test? */
950         CURVE_KN_INTERSECTCORRECTION = 16,              /* correct for width after determing closest midpoint? */
951         CURVE_KN_TRUETANGENTGNORMAL = 32,               /* use tangent normal for geometry? */
952         CURVE_KN_RIBBONS = 64,                                  /* use flat curve ribbons */
953 } CurveFlag;
954
955 typedef struct KernelCurves {
956         int curveflags;
957         int subdivisions;
958
959         float minimum_width;
960         float maximum_width;
961 } KernelCurves;
962
963 typedef struct KernelTables {
964         int beckmann_offset;
965         int pad1, pad2, pad3;
966 } KernelTables;
967
968 typedef struct KernelData {
969         KernelCamera cam;
970         KernelFilm film;
971         KernelBackground background;
972         KernelIntegrator integrator;
973         KernelBVH bvh;
974         KernelCurves curve;
975         KernelTables tables;
976 } KernelData;
977
978 #ifdef __KERNEL_DEBUG__
979 typedef ccl_addr_space struct DebugData {
980         // Total number of BVH node traversal steps and primitives intersections
981         // for the camera rays.
982         int num_bvh_traversal_steps;
983 } DebugData;
984 #endif
985
986 /* Declarations required for split kernel */
987
988 /* Macro for queues */
989 /* Value marking queue's empty slot */
990 #define QUEUE_EMPTY_SLOT -1
991
992 /*
993 * Queue 1 - Active rays
994 * Queue 2 - Background queue
995 * Queue 3 - Shadow ray cast kernel - AO
996 * Queeu 4 - Shadow ray cast kernel - direct lighting
997 */
998 #define NUM_QUEUES 4
999
1000 /* Queue names */
1001 enum QueueNumber {
1002         QUEUE_ACTIVE_AND_REGENERATED_RAYS,         /* All active rays and regenerated rays are enqueued here */
1003         QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS,      /* All
1004                                                     * 1.Background-hit rays,
1005                                                     * 2.Rays that has exited path-iteration but needs to update output buffer
1006                                                     * 3.Rays to be regenerated
1007                                                     * are enqueued here */
1008         QUEUE_SHADOW_RAY_CAST_AO_RAYS,             /* All rays for which a shadow ray should be cast to determine radiance
1009                                                       contribution for AO are enqueued here */
1010         QUEUE_SHADOW_RAY_CAST_DL_RAYS,             /* All rays for which a shadow ray should be cast to determine radiance
1011                                                       contributuin for direct lighting are enqueued here */
1012 };
1013
1014 /* We use RAY_STATE_MASK to get ray_state (enums 0 to 5) */
1015 #define RAY_STATE_MASK 0x007
1016 #define RAY_FLAG_MASK 0x0F8
1017 enum RayState {
1018         RAY_ACTIVE = 0,             // Denotes ray is actively involved in path-iteration
1019         RAY_INACTIVE = 1,           // Denotes ray has completed processing all samples and is inactive
1020         RAY_UPDATE_BUFFER = 2,      // Denoted ray has exited path-iteration and needs to update output buffer
1021         RAY_HIT_BACKGROUND = 3,     // Donotes ray has hit background
1022         RAY_TO_REGENERATE = 4,      // Denotes ray has to be regenerated
1023         RAY_REGENERATED = 5,        // Denotes ray has been regenerated
1024         RAY_SKIP_DL = 6,            // Denotes ray should skip direct lighting
1025         RAY_SHADOW_RAY_CAST_AO = 16, // Flag's ray has to execute shadow blocked function in AO part
1026         RAY_SHADOW_RAY_CAST_DL = 32 // Flag's ray has to execute shadow blocked function in direct lighting part
1027 };
1028
1029 #define ASSIGN_RAY_STATE(ray_state, ray_index, state) (ray_state[ray_index] = ((ray_state[ray_index] & RAY_FLAG_MASK) | state))
1030 #define IS_STATE(ray_state, ray_index, state) ((ray_state[ray_index] & RAY_STATE_MASK) == state)
1031 #define ADD_RAY_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] = (ray_state[ray_index] | flag))
1032 #define REMOVE_RAY_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] = (ray_state[ray_index] & (~flag)))
1033 #define IS_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] & flag)
1034
1035 CCL_NAMESPACE_END
1036
1037 #endif /*  __KERNEL_TYPES_H__ */
1038