path: root/intern/cycles/kernel/kernel_types.h

/*
 * Copyright 2011-2013 Blender Foundation
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

#ifndef __KERNEL_TYPES_H__
#define __KERNEL_TYPES_H__

#include "kernel_math.h"
#include "svm/svm_types.h"

#ifndef __KERNEL_GPU__
#define __KERNEL_CPU__
#endif

/* TODO(sergey): This is only to make it possible to include this header
 * from outside of the kernel. but this could be done somewhat cleaner?
 */
#ifndef ccl_addr_space
#define ccl_addr_space
#endif

CCL_NAMESPACE_BEGIN

/* constants */
#define OBJECT_SIZE 		11
#define OBJECT_VECTOR_SIZE	6
#define LIGHT_SIZE			5
#define FILTER_TABLE_SIZE	256
#define RAMP_TABLE_SIZE		256
#define PARTICLE_SIZE 		5
#define TIME_INVALID		FLT_MAX

#define BSSRDF_MIN_RADIUS			1e-8f
#define BSSRDF_MAX_HITS				4

#define BECKMANN_TABLE_SIZE		256

#define TEX_NUM_FLOAT_IMAGES	5

#define SHADER_NONE				(~0)
#define OBJECT_NONE				(~0)
#define PRIM_NONE				(~0)
#define LAMP_NONE				(~0)

#define VOLUME_STACK_SIZE		16

/* device capabilities */
#ifdef __KERNEL_CPU__
#ifdef __KERNEL_SSE2__
#  define __QBVH__
#endif
#define __KERNEL_SHADING__
#define __KERNEL_ADV_SHADING__
#define __BRANCHED_PATH__
#ifdef WITH_OSL
#define __OSL__
#endif
#define __SUBSURFACE__
#define __CMJ__
#define __VOLUME__
#define __VOLUME_DECOUPLED__
#define __VOLUME_SCATTER__
#define __SHADOW_RECORD_ALL__
#define __VOLUME_RECORD_ALL__
#define __CAMERA_RAY_NODES__
#endif

#ifdef __KERNEL_CUDA__
#define __KERNEL_SHADING__
#define __KERNEL_ADV_SHADING__
#define __BRANCHED_PATH__
#define __VOLUME__
#define __VOLUME_SCATTER__

/* Experimental on GPU */
#ifdef __KERNEL_EXPERIMENTAL__
#define __SUBSURFACE__
#define __CMJ__
#endif

#endif

#ifdef __KERNEL_OPENCL__

/* keep __KERNEL_ADV_SHADING__ in sync with opencl_kernel_use_advanced_shading! */

#ifdef __KERNEL_OPENCL_NVIDIA__
#  define __KERNEL_SHADING__
#  define __KERNEL_ADV_SHADING__
#  ifdef __KERNEL_EXPERIMENTAL__
#    define __CMJ__
#  endif
#endif

#ifdef __KERNEL_OPENCL_APPLE__
#  define __KERNEL_SHADING__
#  define __KERNEL_ADV_SHADING__
/* TODO(sergey): Currently experimental section is ignored here,
 * this is because megakernel in device_opencl does not support
 * custom cflags depending on the scene features.
 */
#  ifdef __KERNEL_EXPERIMENTAL__
#    define __CMJ__
#  endif
#endif

#ifdef __KERNEL_OPENCL_AMD__
#  define __CL_USE_NATIVE__
#  define __KERNEL_SHADING__
#  define __MULTI_CLOSURE__
#  define __PASSES__
#  define __BACKGROUND_MIS__
#  define __LAMP_MIS__
#  define __AO__
#  define __CAMERA_MOTION__
#  define __OBJECT_MOTION__
#  define __HAIR__
#  ifdef __KERNEL_EXPERIMENTAL__
#    define __TRANSPARENT_SHADOWS__
#  endif
#endif

#ifdef __KERNEL_OPENCL_INTEL_CPU__
#  define __CL_USE_NATIVE__
#  define __KERNEL_SHADING__
#  define __KERNEL_ADV_SHADING__
#  ifdef __KERNEL_EXPERIMENTAL__
#    define __CMJ__
#  endif
#endif

#endif // __KERNEL_OPENCL__

/* kernel features */
#define __SOBOL__
#define __INSTANCING__
#define __DPDU__
#define __UV__
#define __BACKGROUND__
#define __CAUSTICS_TRICKS__
#define __VISIBILITY_FLAG__
#define __RAY_DIFFERENTIALS__
#define __CAMERA_CLIPPING__
#define __INTERSECTION_REFINE__
#define __CLAMP_SAMPLE__

#ifdef __KERNEL_SHADING__
#define __SVM__
#define __EMISSION__
#define __TEXTURES__
#define __EXTRA_NODES__
#define __HOLDOUT__
#endif

#ifdef __KERNEL_ADV_SHADING__
#define __MULTI_CLOSURE__
#define __TRANSPARENT_SHADOWS__
#define __PASSES__
#define __BACKGROUND_MIS__
#define __LAMP_MIS__
#define __AO__
#define __CAMERA_MOTION__
#define __OBJECT_MOTION__
#define __HAIR__
#endif

#ifdef WITH_CYCLES_DEBUG
#  define __KERNEL_DEBUG__
#endif

/* Scene-based selective featrues compilation/ */
#ifdef __NO_CAMERA_MOTION__
#  undef __CAMERA_MOTION__
#endif
#ifdef __NO_OBJECT_MOTION__
#  undef __OBJECT_MOTION__
#endif
#ifdef __NO_HAIR__
#  undef __HAIR__
#endif

/* Random Numbers */

typedef uint RNG;

/* Shader Evaluation */

typedef enum ShaderEvalType {
	SHADER_EVAL_DISPLACE,
	SHADER_EVAL_BACKGROUND,
	/* bake types */
	SHADER_EVAL_BAKE, /* no real shade, it's used in the code to
	                   * differentiate the type of shader eval from the above
	                   */
	/* data passes */
	SHADER_EVAL_NORMAL,
	SHADER_EVAL_UV,
	SHADER_EVAL_DIFFUSE_COLOR,
	SHADER_EVAL_GLOSSY_COLOR,
	SHADER_EVAL_TRANSMISSION_COLOR,
	SHADER_EVAL_SUBSURFACE_COLOR,
	SHADER_EVAL_EMISSION,

	/* light passes */
	SHADER_EVAL_AO,
	SHADER_EVAL_COMBINED,
	SHADER_EVAL_SHADOW,
	SHADER_EVAL_DIFFUSE_DIRECT,
	SHADER_EVAL_GLOSSY_DIRECT,
	SHADER_EVAL_TRANSMISSION_DIRECT,
	SHADER_EVAL_SUBSURFACE_DIRECT,
	SHADER_EVAL_DIFFUSE_INDIRECT,
	SHADER_EVAL_GLOSSY_INDIRECT,
	SHADER_EVAL_TRANSMISSION_INDIRECT,
	SHADER_EVAL_SUBSURFACE_INDIRECT,

	/* extra */
	SHADER_EVAL_ENVIRONMENT,
} ShaderEvalType;

/* Path Tracing
 * note we need to keep the u/v pairs at even values */

enum PathTraceDimension {
	PRNG_FILTER_U = 0,
	PRNG_FILTER_V = 1,
	PRNG_LENS_U = 2,
	PRNG_LENS_V = 3,
#ifdef __CAMERA_MOTION__
	PRNG_TIME = 4,
	PRNG_UNUSED_0 = 5,
	PRNG_UNUSED_1 = 6,	/* for some reason (6, 7) is a bad sobol pattern */
	PRNG_UNUSED_2 = 7,  /* with a low number of samples (< 64) */
#endif
	PRNG_BASE_NUM = 8,

	PRNG_BSDF_U = 0,
	PRNG_BSDF_V = 1,
	PRNG_BSDF = 2,
	PRNG_LIGHT = 3,
	PRNG_LIGHT_U = 4,
	PRNG_LIGHT_V = 5,
	PRNG_UNUSED_3 = 6,
	PRNG_TERMINATE = 7,

#ifdef __VOLUME__
	PRNG_PHASE_U = 8,
	PRNG_PHASE_V = 9,
	PRNG_PHASE = 10,
	PRNG_SCATTER_DISTANCE = 11,
#endif

	PRNG_BOUNCE_NUM = 12,
};

enum SamplingPattern {
	SAMPLING_PATTERN_SOBOL = 0,
	SAMPLING_PATTERN_CMJ = 1
};

/* these flags values correspond to raytypes in osl.cpp, so keep them in sync!
 *
 * for ray visibility tests in BVH traversal, the upper 20 bits are used for
 * layer visibility tests. */

enum PathRayFlag {
	PATH_RAY_CAMERA = 1,
	PATH_RAY_REFLECT = 2,
	PATH_RAY_TRANSMIT = 4,
	PATH_RAY_DIFFUSE = 8,
	PATH_RAY_GLOSSY = 16,
	PATH_RAY_SINGULAR = 32,
	PATH_RAY_TRANSPARENT = 64,

	PATH_RAY_SHADOW_OPAQUE = 128,
	PATH_RAY_SHADOW_TRANSPARENT = 256,
	PATH_RAY_SHADOW = (PATH_RAY_SHADOW_OPAQUE|PATH_RAY_SHADOW_TRANSPARENT),

	PATH_RAY_CURVE = 512, /* visibility flag to define curve segments */
	PATH_RAY_VOLUME_SCATTER = 1024, /* volume scattering */

	/* note that these can use maximum 12 bits, the other are for layers */
	PATH_RAY_ALL_VISIBILITY = (1|2|4|8|16|32|64|128|256|512|1024),

	PATH_RAY_MIS_SKIP = 2048,
	PATH_RAY_DIFFUSE_ANCESTOR = 4096,
	PATH_RAY_SINGLE_PASS_DONE = 8192,

	/* we need layer member flags to be the 20 upper bits */
	PATH_RAY_LAYER_SHIFT = (32-20)
};

/* Closure Label */

typedef enum ClosureLabel {
	LABEL_NONE = 0,
	LABEL_TRANSMIT = 1,
	LABEL_REFLECT = 2,
	LABEL_DIFFUSE = 4,
	LABEL_GLOSSY = 8,
	LABEL_SINGULAR = 16,
	LABEL_TRANSPARENT = 32,
	LABEL_VOLUME_SCATTER = 64,
} ClosureLabel;

/* Render Passes */

typedef enum PassType {
	PASS_NONE = 0,
	PASS_COMBINED = (1 << 0),
	PASS_DEPTH = (1 << 1),
	PASS_NORMAL = (1 << 2),
	PASS_UV = (1 << 3),
	PASS_OBJECT_ID = (1 << 4),
	PASS_MATERIAL_ID = (1 << 5),
	PASS_DIFFUSE_COLOR = (1 << 6),
	PASS_GLOSSY_COLOR = (1 << 7),
	PASS_TRANSMISSION_COLOR = (1 << 8),
	PASS_DIFFUSE_INDIRECT = (1 << 9),
	PASS_GLOSSY_INDIRECT = (1 << 10),
	PASS_TRANSMISSION_INDIRECT = (1 << 11),
	PASS_DIFFUSE_DIRECT = (1 << 12),
	PASS_GLOSSY_DIRECT = (1 << 13),
	PASS_TRANSMISSION_DIRECT = (1 << 14),
	PASS_EMISSION = (1 << 15),
	PASS_BACKGROUND = (1 << 16),
	PASS_AO = (1 << 17),
	PASS_SHADOW = (1 << 18),
	PASS_MOTION = (1 << 19),
	PASS_MOTION_WEIGHT = (1 << 20),
	PASS_MIST = (1 << 21),
	PASS_SUBSURFACE_DIRECT = (1 << 22),
	PASS_SUBSURFACE_INDIRECT = (1 << 23),
	PASS_SUBSURFACE_COLOR = (1 << 24),
	PASS_LIGHT = (1 << 25), /* no real pass, used to force use_light_pass */
#ifdef __KERNEL_DEBUG__
	PASS_BVH_TRAVERSAL_STEPS = (1 << 26),
	PASS_BVH_TRAVERSED_INSTANCES = (1 << 27),
	PASS_RAY_BOUNCES = (1 << 28),
#endif
} PassType;

#define PASS_ALL (~0)

#ifdef __PASSES__

typedef ccl_addr_space struct PathRadiance {
	int use_light_pass;

	float3 emission;
	float3 background;
	float3 ao;

	float3 indirect;
	float3 direct_throughput;
	float3 direct_emission;

	float3 color_diffuse;
	float3 color_glossy;
	float3 color_transmission;
	float3 color_subsurface;
	float3 color_scatter;

	float3 direct_diffuse;
	float3 direct_glossy;
	float3 direct_transmission;
	float3 direct_subsurface;
	float3 direct_scatter;

	float3 indirect_diffuse;
	float3 indirect_glossy;
	float3 indirect_transmission;
	float3 indirect_subsurface;
	float3 indirect_scatter;

	float3 path_diffuse;
	float3 path_glossy;
	float3 path_transmission;
	float3 path_subsurface;
	float3 path_scatter;

	float4 shadow;
	float mist;
} PathRadiance;

typedef struct BsdfEval {
	int use_light_pass;

	float3 diffuse;
	float3 glossy;
	float3 transmission;
	float3 transparent;
	float3 subsurface;
	float3 scatter;
} BsdfEval;

#else

typedef ccl_addr_space float3 PathRadiance;
typedef float3 BsdfEval;

#endif

/* Shader Flag */

typedef enum ShaderFlag {
	SHADER_SMOOTH_NORMAL = (1 << 31),
	SHADER_CAST_SHADOW = (1 << 30),
	SHADER_AREA_LIGHT = (1 << 29),
	SHADER_USE_MIS = (1 << 28),
	SHADER_EXCLUDE_DIFFUSE = (1 << 27),
	SHADER_EXCLUDE_GLOSSY = (1 << 26),
	SHADER_EXCLUDE_TRANSMIT = (1 << 25),
	SHADER_EXCLUDE_CAMERA = (1 << 24),
	SHADER_EXCLUDE_SCATTER = (1 << 23),
	SHADER_EXCLUDE_ANY = (SHADER_EXCLUDE_DIFFUSE|SHADER_EXCLUDE_GLOSSY|SHADER_EXCLUDE_TRANSMIT|SHADER_EXCLUDE_CAMERA|SHADER_EXCLUDE_SCATTER),

	SHADER_MASK = ~(SHADER_SMOOTH_NORMAL|SHADER_CAST_SHADOW|SHADER_AREA_LIGHT|SHADER_USE_MIS|SHADER_EXCLUDE_ANY)
} ShaderFlag;

/* Light Type */

typedef enum LightType {
	LIGHT_POINT,
	LIGHT_DISTANT,
	LIGHT_BACKGROUND,
	LIGHT_AREA,
	LIGHT_SPOT,
	LIGHT_TRIANGLE
} LightType;

/* Camera Type */

enum CameraType {
	CAMERA_PERSPECTIVE,
	CAMERA_ORTHOGRAPHIC,
	CAMERA_PANORAMA
};

/* Panorama Type */

enum PanoramaType {
	PANORAMA_EQUIRECTANGULAR,
	PANORAMA_MIRRORBALL,
	PANORAMA_FISHEYE_EQUIDISTANT,
	PANORAMA_FISHEYE_EQUISOLID
};

/* Differential */

typedef struct differential3 {
	float3 dx;
	float3 dy;
} differential3;

typedef struct differential {
	float dx;
	float dy;
} differential;

/* Ray */

typedef struct Ray {
/* TODO(sergey): This is only needed because current AMD
 * compiler has hard time building the kernel with this
 * reshuffle. And at the same time reshuffle will cause
 * less optimal CPU code in certain places.
 *
 * We'll get rid of this nasty exception once AMD compiler
 * is fixed.
 */
#ifndef __KERNEL_OPENCL_AMD__
	float3 P;		/* origin */
	float3 D;		/* direction */

	float t;		/* length of the ray */
	float time;		/* time (for motion blur) */
#else
	float t;		/* length of the ray */
	float time;		/* time (for motion blur) */
	float3 P;		/* origin */
	float3 D;		/* direction */
#endif

#ifdef __RAY_DIFFERENTIALS__
	differential3 dP;
	differential3 dD;
#endif
} Ray;

/* Intersection */

typedef ccl_addr_space struct Intersection {
	float t, u, v;
	int prim;
	int object;
	int type;

#ifdef __KERNEL_DEBUG__
	int num_traversal_steps;
	int num_traversed_instances;
#endif
} Intersection;

/* Primitives */

typedef enum PrimitiveType {
	PRIMITIVE_NONE = 0,
	PRIMITIVE_TRIANGLE = 1,
	PRIMITIVE_MOTION_TRIANGLE = 2,
	PRIMITIVE_CURVE = 4,
	PRIMITIVE_MOTION_CURVE = 8,

	PRIMITIVE_ALL_TRIANGLE = (PRIMITIVE_TRIANGLE|PRIMITIVE_MOTION_TRIANGLE),
	PRIMITIVE_ALL_CURVE = (PRIMITIVE_CURVE|PRIMITIVE_MOTION_CURVE),
	PRIMITIVE_ALL_MOTION = (PRIMITIVE_MOTION_TRIANGLE|PRIMITIVE_MOTION_CURVE),
	PRIMITIVE_ALL = (PRIMITIVE_ALL_TRIANGLE|PRIMITIVE_ALL_CURVE),

	/* Total number of different primitives.
	 * NOTE: This is an actual value, not a bitflag.
	 */
	PRIMITIVE_NUM_TOTAL = 4,
} PrimitiveType;

#define PRIMITIVE_PACK_SEGMENT(type, segment) ((segment << 16) | type)
#define PRIMITIVE_UNPACK_SEGMENT(type) (type >> 16)

/* Attributes */

#define ATTR_PRIM_TYPES		2
#define ATTR_PRIM_CURVE		1

typedef enum AttributeElement {
	ATTR_ELEMENT_NONE,
	ATTR_ELEMENT_OBJECT,
	ATTR_ELEMENT_MESH,
	ATTR_ELEMENT_FACE,
	ATTR_ELEMENT_VERTEX,
	ATTR_ELEMENT_VERTEX_MOTION,
	ATTR_ELEMENT_CORNER,
	ATTR_ELEMENT_CORNER_BYTE,
	ATTR_ELEMENT_CURVE,
	ATTR_ELEMENT_CURVE_KEY,
	ATTR_ELEMENT_CURVE_KEY_MOTION,
	ATTR_ELEMENT_VOXEL
} AttributeElement;

typedef enum AttributeStandard {
	ATTR_STD_NONE = 0,
	ATTR_STD_VERTEX_NORMAL,
	ATTR_STD_FACE_NORMAL,
	ATTR_STD_UV,
	ATTR_STD_UV_TANGENT,
	ATTR_STD_UV_TANGENT_SIGN,
	ATTR_STD_GENERATED,
	ATTR_STD_GENERATED_TRANSFORM,
	ATTR_STD_POSITION_UNDEFORMED,
	ATTR_STD_POSITION_UNDISPLACED,
	ATTR_STD_MOTION_VERTEX_POSITION,
	ATTR_STD_MOTION_VERTEX_NORMAL,
	ATTR_STD_PARTICLE,
	ATTR_STD_CURVE_INTERCEPT,
	ATTR_STD_PTEX_FACE_ID,
	ATTR_STD_PTEX_UV,
	ATTR_STD_VOLUME_DENSITY,
	ATTR_STD_VOLUME_COLOR,
	ATTR_STD_VOLUME_FLAME,
	ATTR_STD_VOLUME_HEAT,
	ATTR_STD_VOLUME_VELOCITY,
	ATTR_STD_POINTINESS,
	ATTR_STD_NUM,

	ATTR_STD_NOT_FOUND = ~0
} AttributeStandard;

/* Closure data */

#ifdef __MULTI_CLOSURE__
#  ifndef __MAX_CLOSURE__
#     define MAX_CLOSURE 64
#  else
#    define MAX_CLOSURE __MAX_CLOSURE__
#  endif
#else
#define MAX_CLOSURE 1
#endif

/* This struct is to be 16 bytes aligned, we also keep some extra precautions:
 * - All the float3 members are in the beginning of the struct, so compiler
 *   does not put own padding trying to align this members.
 * - We make sure OSL pointer is also 16 bytes aligned.
 */
typedef ccl_addr_space struct ShaderClosure {
	float3 weight;
	float3 N;
	float3 T;

	ClosureType type;
	float sample_weight;
	float data0;
	float data1;
	float data2;
	int pad1, pad2, pad3;

#ifdef __OSL__
	void *prim, *pad4;
#endif
} ShaderClosure;

/* Shader Context
 *
 * For OSL we recycle a fixed number of contexts for speed */

typedef enum ShaderContext {
	SHADER_CONTEXT_MAIN = 0,
	SHADER_CONTEXT_INDIRECT = 1,
	SHADER_CONTEXT_EMISSION = 2,
	SHADER_CONTEXT_SHADOW = 3,
	SHADER_CONTEXT_SSS = 4,
	SHADER_CONTEXT_VOLUME = 5,
	SHADER_CONTEXT_NUM = 6
} ShaderContext;

/* Shader Data
 *
 * Main shader state at a point on the surface or in a volume. All coordinates
 * are in world space. */

enum ShaderDataFlag {
	/* runtime flags */
	SD_BACKFACING     = (1 << 0),   /* backside of surface? */
	SD_EMISSION       = (1 << 1),   /* have emissive closure? */
	SD_BSDF           = (1 << 2),   /* have bsdf closure? */
	SD_BSDF_HAS_EVAL  = (1 << 3),   /* have non-singular bsdf closure? */
	SD_BSSRDF         = (1 << 4),   /* have bssrdf */
	SD_HOLDOUT        = (1 << 5),   /* have holdout closure? */
	SD_ABSORPTION     = (1 << 6),   /* have volume absorption closure? */
	SD_SCATTER        = (1 << 7),   /* have volume phase closure? */
	SD_AO             = (1 << 8),   /* have ao closure? */
	SD_TRANSPARENT    = (1 << 9),  /* have transparent closure? */

	SD_CLOSURE_FLAGS = (SD_EMISSION|SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSSRDF|
	                    SD_HOLDOUT|SD_ABSORPTION|SD_SCATTER|SD_AO),

	/* shader flags */
	SD_USE_MIS                = (1 << 10),  /* direct light sample */
	SD_HAS_TRANSPARENT_SHADOW = (1 << 11),  /* has transparent shadow */
	SD_HAS_VOLUME             = (1 << 12),  /* has volume shader */
	SD_HAS_ONLY_VOLUME        = (1 << 13),  /* has only volume shader, no surface */
	SD_HETEROGENEOUS_VOLUME   = (1 << 14),  /* has heterogeneous volume */
	SD_HAS_BSSRDF_BUMP        = (1 << 15),  /* bssrdf normal uses bump */
	SD_VOLUME_EQUIANGULAR     = (1 << 16),  /* use equiangular sampling */
	SD_VOLUME_MIS             = (1 << 17),  /* use multiple importance sampling */
	SD_VOLUME_CUBIC           = (1 << 18),  /* use cubic interpolation for voxels */
	SD_HAS_BUMP               = (1 << 19),  /* has data connected to the displacement input */

	SD_SHADER_FLAGS = (SD_USE_MIS|SD_HAS_TRANSPARENT_SHADOW|SD_HAS_VOLUME|
	                   SD_HAS_ONLY_VOLUME|SD_HETEROGENEOUS_VOLUME|
	                   SD_HAS_BSSRDF_BUMP|SD_VOLUME_EQUIANGULAR|SD_VOLUME_MIS|
	                   SD_VOLUME_CUBIC|SD_HAS_BUMP),

	/* object flags */
	SD_HOLDOUT_MASK             = (1 << 20),  /* holdout for camera rays */
	SD_OBJECT_MOTION            = (1 << 21),  /* has object motion blur */
	SD_TRANSFORM_APPLIED        = (1 << 22),  /* vertices have transform applied */
	SD_NEGATIVE_SCALE_APPLIED   = (1 << 23),  /* vertices have negative scale applied */
	SD_OBJECT_HAS_VOLUME        = (1 << 24),  /* object has a volume shader */
	SD_OBJECT_INTERSECTS_VOLUME = (1 << 25),  /* object intersects AABB of an object with volume shader */
	SD_OBJECT_HAS_VERTEX_MOTION = (1 << 26),  /* has position for motion vertices */

	SD_OBJECT_FLAGS = (SD_HOLDOUT_MASK|SD_OBJECT_MOTION|SD_TRANSFORM_APPLIED|
	                   SD_NEGATIVE_SCALE_APPLIED|SD_OBJECT_HAS_VOLUME|
	                   SD_OBJECT_INTERSECTS_VOLUME)
};

struct KernelGlobals;

#ifdef __SPLIT_KERNEL__
#define SD_VAR(type, what) ccl_global type *what;
#define SD_CLOSURE_VAR(type, what, max_closure) type *what;
#define TIDX (get_global_id(1) * get_global_size(0) + get_global_id(0))
#define ccl_fetch(s, t) (s->t[TIDX])
#define ccl_fetch_array(s, t, index) (&s->t[TIDX * MAX_CLOSURE + index])
#else
#define SD_VAR(type, what) type what;
#define SD_CLOSURE_VAR(type, what, max_closure) type what[max_closure];
#define ccl_fetch(s, t) (s->t)
#define ccl_fetch_array(s, t, index) (&s->t[index])
#endif

typedef ccl_addr_space struct ShaderData {

#include "kernel_shaderdata_vars.h"

} ShaderData;

/* Path State */

#ifdef __VOLUME__
typedef struct VolumeStack {
	int object;
	int shader;
} VolumeStack;
#endif

typedef struct PathState {
	/* see enum PathRayFlag */
	int flag;          

	/* random number generator state */
	int rng_offset;    		/* dimension offset */
	int sample;        		/* path sample number */
	int num_samples;		/* total number of times this path will be sampled */

	/* bounce counting */
	int bounce;
	int diffuse_bounce;
	int glossy_bounce;
	int transmission_bounce;
	int transparent_bounce;

	/* multiple importance sampling */
	float min_ray_pdf; /* smallest bounce pdf over entire path up to now */
	float ray_pdf;     /* last bounce pdf */
#ifdef __LAMP_MIS__
	float ray_t;       /* accumulated distance through transparent surfaces */
#endif

	/* volume rendering */
#ifdef __VOLUME__
	int volume_bounce;
	RNG rng_congruential;
	VolumeStack volume_stack[VOLUME_STACK_SIZE];
#endif
} PathState;

/* Constant Kernel Data
 *
 * These structs are passed from CPU to various devices, and the struct layout
 * must match exactly. Structs are padded to ensure 16 byte alignment, and we
 * do not use float3 because its size may not be the same on all devices. */

typedef struct KernelCamera {
	/* type */
	int type;

	/* panorama */
	int panorama_type;
	float fisheye_fov;
	float fisheye_lens;
	float4 equirectangular_range;

	/* matrices */
	Transform cameratoworld;
	Transform rastertocamera;

	/* differentials */
	float4 dx;
	float4 dy;

	/* depth of field */
	float aperturesize;
	float blades;
	float bladesrotation;
	float focaldistance;

	/* motion blur */
	float shuttertime;
	int have_motion, have_perspective_motion;

	/* clipping */
	float nearclip;
	float cliplength;

	/* sensor size */
	float sensorwidth;
	float sensorheight;

	/* render size */
	float width, height;
	int resolution;

	/* anamorphic lens bokeh */
	float inv_aperture_ratio;

	int is_inside_volume;
	int shader;

	float focal_length;
	float pad[2];

	/* more matrices */
	Transform screentoworld;
	Transform rastertoworld;
	/* work around cuda sm 2.0 crash, this seems to
	 * cross some limit in combination with motion 
	 * Transform ndctoworld; */
	Transform worldtoscreen;
	Transform worldtoraster;
	Transform worldtondc;
	Transform worldtocamera;

	MotionTransform motion;

	/* Denotes changes in the projective matrix, namely in rastertocamera.
	 * Used for camera zoom motion blur,
	 */
	PerspectiveMotionTransform perspective_motion;
} KernelCamera;

typedef struct KernelFilm {
	float exposure;
	int pass_flag;
	int pass_stride;
	int use_light_pass;

	int pass_combined;
	int pass_depth;
	int pass_normal;
	int pass_motion;

	int pass_motion_weight;
	int pass_uv;
	int pass_object_id;
	int pass_material_id;

	int pass_diffuse_color;
	int pass_glossy_color;
	int pass_transmission_color;
	int pass_subsurface_color;
	
	int pass_diffuse_indirect;
	int pass_glossy_indirect;
	int pass_transmission_indirect;
	int pass_subsurface_indirect;
	
	int pass_diffuse_direct;
	int pass_glossy_direct;
	int pass_transmission_direct;
	int pass_subsurface_direct;
	
	int pass_emission;
	int pass_background;
	int pass_ao;
	float pass_alpha_threshold;

	int pass_shadow;
	float pass_shadow_scale;
	int filter_table_offset;
	int pass_pad2;

	int pass_mist;
	float mist_start;
	float mist_inv_depth;
	float mist_falloff;

#ifdef __KERNEL_DEBUG__
	int pass_bvh_traversal_steps;
	int pass_bvh_traversed_instances;
	int pass_ray_bounces;
	int pass_pad3;
#endif
} KernelFilm;

typedef struct KernelBackground {
	/* only shader index */
	int surface_shader;
	int volume_shader;
	int transparent;
	int pad;

	/* ambient occlusion */
	float ao_factor;
	float ao_distance;
	float ao_pad1, ao_pad2;
} KernelBackground;

typedef struct KernelIntegrator {
	/* emission */
	int use_direct_light;
	int use_ambient_occlusion;
	int num_distribution;
	int num_all_lights;
	float pdf_triangles;
	float pdf_lights;
	float inv_pdf_lights;
	int pdf_background_res;

	/* light portals */
	float portal_pdf;
	int num_portals;
	int portal_offset;

	/* bounces */
	int min_bounce;
	int max_bounce;

	int max_diffuse_bounce;
	int max_glossy_bounce;
	int max_transmission_bounce;
	int max_volume_bounce;

	/* transparent */
	int transparent_min_bounce;
	int transparent_max_bounce;
	int transparent_shadows;

	/* caustics */
	int caustics_reflective;
	int caustics_refractive;
	float filter_glossy;

	/* seed */
	int seed;

	/* render layer */
	int layer_flag;

	/* clamp */
	float sample_clamp_direct;
	float sample_clamp_indirect;

	/* branched path */
	int branched;
	int diffuse_samples;
	int glossy_samples;
	int transmission_samples;
	int ao_samples;
	int mesh_light_samples;
	int subsurface_samples;
	int sample_all_lights_direct;
	int sample_all_lights_indirect;

	/* mis */
	int use_lamp_mis;

	/* sampler */
	int sampling_pattern;
	int aa_samples;

	/* volume render */
	int use_volumes;
	int volume_max_steps;
	float volume_step_size;
	int volume_samples;

	int pad;
} KernelIntegrator;

typedef struct KernelBVH {
	/* root node */
	int root;
	int attributes_map_stride;
	int have_motion;
	int have_curves;
	int have_instancing;
	int use_qbvh;
	int pad1, pad2;
} KernelBVH;

typedef enum CurveFlag {
	/* runtime flags */
	CURVE_KN_BACKFACING = 1,				/* backside of cylinder? */
	CURVE_KN_ENCLOSEFILTER = 2,				/* don't consider strands surrounding start point? */
	CURVE_KN_INTERPOLATE = 4,				/* render as a curve? */
	CURVE_KN_ACCURATE = 8,					/* use accurate intersections test? */
	CURVE_KN_INTERSECTCORRECTION = 16,		/* correct for width after determing closest midpoint? */
	CURVE_KN_TRUETANGENTGNORMAL = 32,		/* use tangent normal for geometry? */
	CURVE_KN_RIBBONS = 64,					/* use flat curve ribbons */
} CurveFlag;

typedef struct KernelCurves {
	int curveflags;
	int subdivisions;

	float minimum_width;
	float maximum_width;
} KernelCurves;

typedef struct KernelTables {
	int beckmann_offset;
	int pad1, pad2, pad3;
} KernelTables;

typedef struct KernelData {
	KernelCamera cam;
	KernelFilm film;
	KernelBackground background;
	KernelIntegrator integrator;
	KernelBVH bvh;
	KernelCurves curve;
	KernelTables tables;
} KernelData;

typedef struct CameraData {
	float2 raster;
	float2 lens;
	Ray ray;
	int shader;
} CameraData;

#ifdef __KERNEL_DEBUG__
typedef ccl_addr_space struct DebugData {
	// Total number of BVH node traversal steps and primitives intersections
	// for the camera rays.
	int num_bvh_traversal_steps;
	int num_bvh_traversed_instances;
	int num_ray_bounces;
} DebugData;
#endif

/* Declarations required for split kernel */

/* Macro for queues */
/* Value marking queue's empty slot */
#define QUEUE_EMPTY_SLOT -1

/*
* Queue 1 - Active rays
* Queue 2 - Background queue
* Queue 3 - Shadow ray cast kernel - AO
* Queeu 4 - Shadow ray cast kernel - direct lighting
*/
#define NUM_QUEUES 4

/* Queue names */
enum QueueNumber {
	QUEUE_ACTIVE_AND_REGENERATED_RAYS = 0,     /* All active rays and regenerated rays are enqueued here. */
	QUEUE_HITBG_BUFF_UPDATE_TOREGEN_RAYS = 1,  /* All
	                                            * 1. Background-hit rays,
	                                            * 2. Rays that has exited path-iteration but needs to update output buffer
	                                            * 3. Rays to be regenerated
	                                            * are enqueued here.
	                                            */
	QUEUE_SHADOW_RAY_CAST_AO_RAYS = 2,         /* All rays for which a shadow ray should be cast to determine radiance
	                                            * contribution for AO are enqueued here.
	                                            */
	QUEUE_SHADOW_RAY_CAST_DL_RAYS = 3,         /* All rays for which a shadow ray should be cast to determine radiance
	                                            * contributing for direct lighting are enqueued here.
	                                            */
};

/* We use RAY_STATE_MASK to get ray_state (enums 0 to 5) */
#define RAY_STATE_MASK 0x007
#define RAY_FLAG_MASK 0x0F8
enum RayState {
	RAY_ACTIVE = 0,             // Denotes ray is actively involved in path-iteration
	RAY_INACTIVE = 1,           // Denotes ray has completed processing all samples and is inactive
	RAY_UPDATE_BUFFER = 2,      // Denoted ray has exited path-iteration and needs to update output buffer
	RAY_HIT_BACKGROUND = 3,     // Donotes ray has hit background
	RAY_TO_REGENERATE = 4,      // Denotes ray has to be regenerated
	RAY_REGENERATED = 5,        // Denotes ray has been regenerated
	RAY_SKIP_DL = 6,            // Denotes ray should skip direct lighting
	RAY_SHADOW_RAY_CAST_AO = 16, // Flag's ray has to execute shadow blocked function in AO part
	RAY_SHADOW_RAY_CAST_DL = 32 // Flag's ray has to execute shadow blocked function in direct lighting part
};

#define ASSIGN_RAY_STATE(ray_state, ray_index, state) (ray_state[ray_index] = ((ray_state[ray_index] & RAY_FLAG_MASK) | state))
#define IS_STATE(ray_state, ray_index, state) ((ray_state[ray_index] & RAY_STATE_MASK) == state)
#define ADD_RAY_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] = (ray_state[ray_index] | flag))
#define REMOVE_RAY_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] = (ray_state[ray_index] & (~flag)))
#define IS_FLAG(ray_state, ray_index, flag) (ray_state[ray_index] & flag)

CCL_NAMESPACE_END

#endif /*  __KERNEL_TYPES_H__ */