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