Cycles: Render Passes

Currently supported passes:
* Combined, Z, Normal, Object Index, Material Index, Emission, Environment,
  Diffuse/Glossy/Transmission x Direct/Indirect/Color

Not supported yet:
* UV, Vector, Mist

Only enabled for CPU devices at the moment, will do GPU tweaks tommorrow,
also for environment importance sampling.

Documentation:
http://wiki.blender.org/index.php/Doc:2.6/Manual/Render/Cycles/Passes

16 files changed, 758 insertions, 128 deletions

diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index b6e749f3fcb..59fe9709e61 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -15,6 +15,7 @@ set(SRC
 
 set(SRC_HEADERS
 	kernel.h
+	kernel_accumulate.h
 	kernel_bvh.h
 	kernel_camera.h
 	kernel_compat_cpu.h
@@ -30,6 +31,7 @@ set(SRC_HEADERS
 	kernel_mbvh.h
 	kernel_montecarlo.h
 	kernel_object.h
+	kernel_passes.h
 	kernel_path.h
 	kernel_qbvh.h
 	kernel_random.h
diff --git a/intern/cycles/kernel/kernel.cl b/intern/cycles/kernel/kernel.cl
index 305d81339c2..f98414d4f02 100644
--- a/intern/cycles/kernel/kernel.cl
+++ b/intern/cycles/kernel/kernel.cl
@@ -28,7 +28,7 @@
 
 __kernel void kernel_ocl_path_trace(
 	__constant KernelData *data,
-	__global float4 *buffer,
+	__global float *buffer,
 	__global uint *rng_state,
 
 #define KERNEL_TEX(type, ttype, name) \
@@ -56,7 +56,7 @@ __kernel void kernel_ocl_path_trace(
 __kernel void kernel_ocl_tonemap(
 	__constant KernelData *data,
 	__global uchar4 *rgba,
-	__global float4 *buffer,
+	__global float *buffer,
 
 #define KERNEL_TEX(type, ttype, name) \
 	__global type *name,
diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp
index 9e0d252772b..a93f6172d28 100644
--- a/intern/cycles/kernel/kernel.cpp
+++ b/intern/cycles/kernel/kernel.cpp
@@ -204,14 +204,14 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
 
 /* Path Tracing */
 
-void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
+void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
 	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
 
-void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
 	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
diff --git a/intern/cycles/kernel/kernel.cu b/intern/cycles/kernel/kernel.cu
index 4e585fd563d..fc7992c87f0 100644
--- a/intern/cycles/kernel/kernel.cu
+++ b/intern/cycles/kernel/kernel.cu
@@ -26,7 +26,7 @@
 #include "kernel_path.h"
 #include "kernel_displace.h"
 
-extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
+extern "C" __global__ void kernel_cuda_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
 	int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
@@ -35,7 +35,7 @@ extern "C" __global__ void kernel_cuda_path_trace(float4 *buffer, uint *rng_stat
 		kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
 }
 
-extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float4 *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride)
+extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int resolution, int sx, int sy, int sw, int sh, int offset, int stride)
 {
 	int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
 	int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
diff --git a/intern/cycles/kernel/kernel.h b/intern/cycles/kernel/kernel.h
index df6b5ee92da..26c0bcd6d1a 100644
--- a/intern/cycles/kernel/kernel.h
+++ b/intern/cycles/kernel/kernel.h
@@ -36,17 +36,17 @@ bool kernel_osl_use(KernelGlobals *kg);
 void kernel_const_copy(KernelGlobals *kg, const char *name, void *host, size_t size);
 void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t width, size_t height);
 
-void kernel_cpu_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
+void kernel_cpu_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state,
 	int sample, int x, int y, int offset, int stride);
-void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
+void kernel_cpu_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer,
 	int sample, int resolution, int x, int y, int offset, int stride);
 void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output,
 	int type, int i);
 
 #ifdef WITH_OPTIMIZED_KERNEL
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state,
+void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state,
 	int sample, int x, int y, int offset, int stride);
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer,
+void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer,
 	int sample, int resolution, int x, int y, int offset, int stride);
 void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output,
 	int type, int i);
diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h
new file mode 100644
index 00000000000..e71fad58c96
--- /dev/null
+++ b/intern/cycles/kernel/kernel_accumulate.h
@@ -0,0 +1,283 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ */
+
+CCL_NAMESPACE_BEGIN
+
+/* BSDF Eval
+ *
+ * BSDF evaluation result, split per BSDF type. This is used to accumulate
+ * render passes separately. */
+
+__device_inline void bsdf_eval_init(BsdfEval *eval, ClosureType type, float3 value, int use_light_pass)
+{
+#ifdef __PASSES__
+	eval->use_light_pass = use_light_pass;
+
+	if(eval->use_light_pass) {
+		eval->diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		eval->glossy = make_float3(0.0f, 0.0f, 0.0f);
+		eval->transmission = make_float3(0.0f, 0.0f, 0.0f);
+		eval->transparent = make_float3(0.0f, 0.0f, 0.0f);
+
+		if(type == CLOSURE_BSDF_TRANSPARENT_ID)
+			eval->transparent = value;
+		else if(CLOSURE_IS_BSDF_DIFFUSE(type))
+			eval->diffuse = value;
+		else if(CLOSURE_IS_BSDF_GLOSSY(type))
+			eval->glossy = value;
+		else
+			eval->transmission = value;
+	}
+	else
+		eval->diffuse = value;
+#else
+	*eval = value;
+#endif
+}
+
+__device_inline void bsdf_eval_accum(BsdfEval *eval, ClosureType type, float3 value)
+{
+#ifdef __PASSES__
+	if(eval->use_light_pass) {
+		if(CLOSURE_IS_BSDF_DIFFUSE(type))
+			eval->diffuse += value;
+		else if(CLOSURE_IS_BSDF_GLOSSY(type))
+			eval->glossy += value;
+		else
+			eval->transmission += value;
+
+		/* skipping transparent, this function is used by for eval(), will be zero then */
+	}
+	else
+		eval->diffuse += value;
+#else
+	*eval += value;
+#endif
+}
+
+__device_inline bool bsdf_eval_is_zero(BsdfEval *eval)
+{
+#ifdef __PASSES__
+	if(eval->use_light_pass) {
+		return is_zero(eval->diffuse)
+			&& is_zero(eval->glossy)
+			&& is_zero(eval->transmission)
+			&& is_zero(eval->transparent);
+	}
+	else
+		return is_zero(eval->diffuse);
+#else
+	return is_zero(*eval);
+#endif
+}
+
+__device_inline void bsdf_eval_mul(BsdfEval *eval, float3 value)
+{
+#ifdef __PASSES__
+	if(eval->use_light_pass) {
+		eval->diffuse *= value;
+		eval->glossy *= value;
+		eval->transmission *= value;
+
+		/* skipping transparent, this function is used by for eval(), will be zero then */
+	}
+	else
+		eval->diffuse *= value;
+#else
+	*eval *= value;
+#endif
+}
+
+/* Path Radiance
+ *
+ * We accumulate different render passes separately. After summing at the end
+ * to get the combined result, it should be identical. We definte directly
+ * visible as the first non-transparent hit, while indirectly visible are the
+ * bounces after that. */
+
+__device_inline void path_radiance_init(PathRadiance *L, int use_light_pass)
+{
+	/* clear all */
+#ifdef __PASSES__
+	L->use_light_pass = use_light_pass;
+
+	if(use_light_pass) {
+		L->indirect = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_throughput = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_emission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->color_diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		L->color_glossy = make_float3(0.0f, 0.0f, 0.0f);
+		L->color_transmission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->direct_diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_glossy = make_float3(0.0f, 0.0f, 0.0f);
+		L->direct_transmission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->indirect_diffuse = make_float3(0.0f, 0.0f, 0.0f);
+		L->indirect_glossy = make_float3(0.0f, 0.0f, 0.0f);
+		L->indirect_transmission = make_float3(0.0f, 0.0f, 0.0f);
+
+		L->emission = make_float3(0.0f, 0.0f, 0.0f);
+		L->background = make_float3(0.0f, 0.0f, 0.0f);
+	}
+	else
+		L->emission = make_float3(0.0f, 0.0f, 0.0f);
+#else
+	*L = make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
+
+__device_inline void path_radiance_bsdf_bounce(PathRadiance *L, float3 *throughput,
+	BsdfEval *bsdf_eval, float bsdf_pdf, int bounce, int bsdf_label)
+{
+	float inverse_pdf = 1.0f/bsdf_pdf;
+
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0) {
+			if(bsdf_label & LABEL_TRANSPARENT) {
+				/* transparent bounce before first hit */
+				*throughput *= bsdf_eval->transparent*inverse_pdf;
+			}
+			else {
+				/* first on directly visible surface */
+				float3 value = *throughput*inverse_pdf;
+
+				L->indirect_diffuse = bsdf_eval->diffuse*value;
+				L->indirect_glossy = bsdf_eval->glossy*value;
+				L->indirect_transmission = bsdf_eval->transmission*value;
+
+				*throughput = L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission;
+				
+				L->direct_throughput = *throughput;
+			}
+		}
+		else {
+			/* indirectly visible through BSDF */
+			float3 sum = (bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission + bsdf_eval->transparent)*inverse_pdf;
+			*throughput *= sum;
+		}
+	}
+	else
+		*throughput *= bsdf_eval->diffuse*inverse_pdf;
+#else
+	*throughput *= *bsdf_eval*inverse_pdf;
+#endif
+}
+
+__device_inline void path_radiance_accum_emission(PathRadiance *L, float3 throughput, float3 value, int bounce)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0)
+			L->emission += throughput*value;
+		else if(bounce == 1)
+			L->direct_emission += throughput*value;
+		else
+			L->indirect += throughput*value;
+	}
+	else
+		L->emission += throughput*value;
+#else
+	*L += throughput*value;
+#endif
+}
+
+__device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, int bounce)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0) {
+			/* directly visible lighting */
+			L->direct_diffuse += throughput*bsdf_eval->diffuse;
+			L->direct_glossy += throughput*bsdf_eval->glossy;
+			L->direct_transmission += throughput*bsdf_eval->transmission;
+		}
+		else {
+			/* indirectly visible lighting after BSDF bounce */
+			float3 sum = bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission;
+			L->indirect += throughput*sum;
+		}
+	}
+	else
+		L->emission += throughput*bsdf_eval->diffuse;
+#else
+	*L += throughput*(*bsdf_eval);
+#endif
+}
+
+__device_inline void path_radiance_accum_background(PathRadiance *L, float3 throughput, float3 value, int bounce)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		if(bounce == 0)
+			L->background += throughput*value;
+		else if(bounce == 1)
+			L->direct_emission += throughput*value;
+		else
+			L->indirect += throughput*value;
+	}
+	else
+		L->emission += throughput*value;
+#else
+	*L += throughput*value;
+#endif
+}
+
+__device_inline float3 safe_divide_color(float3 a, float3 b)
+{
+	float x, y, z;
+
+	x = (b.x != 0.0f)? a.x/b.x: 0.0f;
+	y = (b.y != 0.0f)? a.y/b.y: 0.0f;
+	z = (b.z != 0.0f)? a.z/b.z: 0.0f;
+
+	return make_float3(x, y, z);
+}
+
+__device_inline float3 path_radiance_sum(PathRadiance *L)
+{
+#ifdef __PASSES__
+	if(L->use_light_pass) {
+		/* this division is a bit ugly, but means we only have to keep track of
+		   only a single throughput further along the path, here we recover just
+		   the indirect parth that is not influenced by any particular BSDF type */
+		L->direct_emission = safe_divide_color(L->direct_emission, L->direct_throughput);
+		L->direct_diffuse += L->indirect_diffuse*L->direct_emission;
+		L->direct_glossy += L->indirect_glossy*L->direct_emission;
+		L->direct_transmission += L->indirect_transmission*L->direct_emission;
+
+		L->indirect = safe_divide_color(L->indirect, L->direct_throughput);
+		L->indirect_diffuse *= L->indirect;
+		L->indirect_glossy *= L->indirect;
+		L->indirect_transmission *= L->indirect;
+
+		return L->emission + L->background
+			+ L->direct_diffuse + L->direct_glossy + L->direct_transmission
+			+ L->indirect_diffuse + L->indirect_glossy + L->indirect_transmission;
+	}
+	else
+		return L->emission;
+#else
+	return *L;
+#endif
+}
+
+CCL_NAMESPACE_END
+
diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h
index 51698f3a9bd..b3a5b2bfcb4 100644
--- a/intern/cycles/kernel/kernel_emission.h
+++ b/intern/cycles/kernel/kernel_emission.h
@@ -57,7 +57,7 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando,
 }
 
 __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
-	float randt, float rando, float randu, float randv, Ray *ray, float3 *eval)
+	float randt, float rando, float randu, float randv, Ray *ray, BsdfEval *eval)
 {
 	LightSample ls;
 
@@ -83,32 +83,33 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
 		return false;
 
 	/* evaluate closure */
-	*eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D);
+	float3 light_eval = direct_emissive_eval(kg, rando, &ls, randu, randv, -ls.D);
 
-	if(is_zero(*eval))
+	if(is_zero(light_eval))
 		return false;
 
 	/* todo: use visbility flag to skip lights */
 
 	/* evaluate BSDF at shading point */
 	float bsdf_pdf;
-	float3 bsdf_eval = shader_bsdf_eval(kg, sd, ls.D, &bsdf_pdf);
 
-	*eval *= bsdf_eval/pdf;
-
-	if(is_zero(*eval))
-		return false;
+	shader_bsdf_eval(kg, sd, ls.D, eval, &bsdf_pdf);
 
 	if(ls.prim != ~0 || ls.type == LIGHT_BACKGROUND) {
 		/* multiple importance sampling */
 		float mis_weight = power_heuristic(pdf, bsdf_pdf);
-		*eval *= mis_weight;
+		light_eval *= mis_weight;
 	}
 	/* todo: clean up these weights */
 	else if(ls.shader & SHADER_AREA_LIGHT)
-		*eval *= 0.25f; /* area lamp */
+		light_eval *= 0.25f; /* area lamp */
 	else if(ls.t != FLT_MAX)
-		*eval *= 0.25f*M_1_PI_F; /* point lamp */
+		light_eval *= 0.25f*M_1_PI_F; /* point lamp */
+	
+	bsdf_eval_mul(eval, light_eval/pdf);
+
+	if(bsdf_eval_is_zero(eval))
+		return false;
 
 	if(ls.shader & SHADER_CAST_SHADOW) {
 		/* setup ray */
diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h
index cd8acc9647a..d0fb5402291 100644
--- a/intern/cycles/kernel/kernel_film.h
+++ b/intern/cycles/kernel/kernel_film.h
@@ -48,15 +48,22 @@ __device uchar4 film_float_to_byte(float4 color)
 	return result;
 }
 
-__device void kernel_film_tonemap(KernelGlobals *kg, __global uchar4 *rgba, __global float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+__device void kernel_film_tonemap(KernelGlobals *kg,
+	__global uchar4 *rgba, __global float *buffer,
+	int sample, int resolution, int x, int y, int offset, int stride)
 {
+	/* buffer offset */
 	int index = offset + x + y*stride;
-	float4 irradiance = buffer[index];
 
+	rgba += index;
+	buffer += index*kernel_data.film.pass_stride;
+
+	/* map colors */
+	float4 irradiance = *(float4*)buffer;
 	float4 float_result = film_map(kg, irradiance, sample);
 	uchar4 byte_result = film_float_to_byte(float_result);
 
-	rgba[index] = byte_result;
+	*rgba = byte_result;
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_object.h b/intern/cycles/kernel/kernel_object.h
index b4c42d3bacc..318a6fea489 100644
--- a/intern/cycles/kernel/kernel_object.h
+++ b/intern/cycles/kernel/kernel_object.h
@@ -64,5 +64,15 @@ __device_inline float object_surface_area(KernelGlobals *kg, int object)
 	return f.x;
 }
 
+__device_inline float object_pass_id(KernelGlobals *kg, int object)
+{
+	if(object == ~0)
+		return 0.0f;
+
+	int offset = object*OBJECT_SIZE + OBJECT_PROPERTIES;
+	float4 f = kernel_tex_fetch(__objects, offset);
+	return f.y;
+}
+
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/kernel/kernel_optimized.cpp b/intern/cycles/kernel/kernel_optimized.cpp
index 50341021d9d..393686bb203 100644
--- a/intern/cycles/kernel/kernel_optimized.cpp
+++ b/intern/cycles/kernel/kernel_optimized.cpp
@@ -35,14 +35,14 @@ CCL_NAMESPACE_BEGIN
 
 /* Path Tracing */
 
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float4 *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
+void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
 	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
 
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float4 *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
 	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h
new file mode 100644
index 00000000000..0e775812eda
--- /dev/null
+++ b/intern/cycles/kernel/kernel_passes.h
@@ -0,0 +1,146 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ */
+
+CCL_NAMESPACE_BEGIN
+
+__device_inline void kernel_write_pass_float(__global float *buffer, int sample, float value)
+{
+	float *buf = buffer;
+	*buf = (sample == 0)? value: *buf + value;
+}
+
+__device_inline void kernel_write_pass_float3(__global float *buffer, int sample, float3 value)
+{
+	float3 *buf = (float3*)buffer;
+	*buf = (sample == 0)? value: *buf + value;
+}
+
+__device_inline void kernel_write_pass_float4(__global float *buffer, int sample, float4 value)
+{
+	float4 *buf = (float4*)buffer;
+	*buf = (sample == 0)? value: *buf + value;
+}
+
+__device_inline void kernel_clear_passes(__global float *buffer, int sample, int pass_stride)
+{
+#ifdef __PASSES__
+	if(sample == 0 && pass_stride != 4)
+		for(int i = 4; i < pass_stride; i++)
+			buffer[i] = 0.0f;
+#endif
+}
+
+__device void kernel_write_data_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L,
+	ShaderData *sd, int sample, int path_flag, float3 throughput)
+{
+#ifdef __PASSES__
+	if(!(path_flag & PATH_RAY_CAMERA))
+		return;
+
+	int flag = kernel_data.film.pass_flag;
+
+	if(!(flag & PASS_ALL))
+		return;
+	
+	/* todo: add alpha treshold */
+	if(!(path_flag & PATH_RAY_TRANSPARENT)) {
+		if(sample == 0) {
+			if(flag & PASS_DEPTH) {
+				Transform tfm = kernel_data.cam.worldtocamera;
+				float depth = len(transform(&tfm, sd->P));
+
+				kernel_write_pass_float(buffer + kernel_data.film.pass_depth, sample, depth);
+			}
+			if(flag & PASS_OBJECT_ID) {
+				float id = object_pass_id(kg, sd->object);
+				kernel_write_pass_float(buffer + kernel_data.film.pass_object_id, sample, id);
+			}
+			if(flag & PASS_MATERIAL_ID) {
+				float id = shader_pass_id(kg, sd);
+				kernel_write_pass_float(buffer + kernel_data.film.pass_material_id, sample, id);
+			}
+		}
+
+		if(flag & PASS_NORMAL) {
+			float3 normal = sd->N;
+			kernel_write_pass_float3(buffer + kernel_data.film.pass_normal, sample, normal);
+		}
+		if(flag & PASS_UV) {
+			float3 uv = make_float3(0.0f, 0.0f, 0.0f); /* todo: request and lookup */
+			kernel_write_pass_float3(buffer + kernel_data.film.pass_uv, sample, uv);
+		}
+	}
+
+	if(flag & (PASS_DIFFUSE_INDIRECT|PASS_DIFFUSE_COLOR|PASS_DIFFUSE_DIRECT))
+		L->color_diffuse += shader_bsdf_diffuse(kg, sd)*throughput;
+	if(flag & (PASS_GLOSSY_INDIRECT|PASS_GLOSSY_COLOR|PASS_GLOSSY_DIRECT))
+		L->color_glossy += shader_bsdf_glossy(kg, sd)*throughput;
+	if(flag & (PASS_TRANSMISSION_INDIRECT|PASS_TRANSMISSION_COLOR|PASS_TRANSMISSION_DIRECT))
+		L->color_transmission += shader_bsdf_transmission(kg, sd)*throughput;
+#endif
+}
+
+__device void kernel_write_light_passes(KernelGlobals *kg, __global float *buffer, PathRadiance *L, int sample)
+{
+#ifdef __PASSES__
+	int flag = kernel_data.film.pass_flag;
+
+	if(!kernel_data.film.use_light_pass)
+		return;
+	
+	if(flag & PASS_DIFFUSE_INDIRECT) {
+		float3 color = safe_divide_color(L->indirect_diffuse, L->color_diffuse);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_indirect, sample, color);
+	}
+	if(flag & PASS_GLOSSY_INDIRECT) {
+		float3 color = safe_divide_color(L->indirect_glossy, L->color_glossy);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_indirect, sample, color);
+	}
+	if(flag & PASS_TRANSMISSION_INDIRECT) {
+		float3 color = safe_divide_color(L->indirect_transmission, L->color_transmission);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_indirect, sample, color);
+	}
+	if(flag & PASS_DIFFUSE_DIRECT) {
+		float3 color = safe_divide_color(L->direct_diffuse, L->color_diffuse);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_direct, sample, color);
+	}
+	if(flag & PASS_GLOSSY_DIRECT) {
+		float3 color = safe_divide_color(L->direct_glossy, L->color_glossy);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_direct, sample, color);
+	}
+	if(flag & PASS_TRANSMISSION_DIRECT) {
+		float3 color = safe_divide_color(L->direct_transmission, L->color_transmission);
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_direct, sample, color);
+	}
+
+	if(flag & PASS_EMISSION)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_emission, sample, L->emission);
+	if(flag & PASS_BACKGROUND)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_background, sample, L->background);
+
+	if(flag & PASS_DIFFUSE_COLOR)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_diffuse_color, sample, L->color_diffuse);
+	if(flag & PASS_GLOSSY_COLOR)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_glossy_color, sample, L->color_glossy);
+	if(flag & PASS_TRANSMISSION_COLOR)
+		kernel_write_pass_float3(buffer + kernel_data.film.pass_transmission_color, sample, L->color_transmission);
+#endif
+}
+
+CCL_NAMESPACE_END
+
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index c80d2068506..c0bfa320405 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -25,39 +25,16 @@
 #else
 #include "kernel_bvh.h"
 #endif
+#include "kernel_accumulate.h"
 #include "kernel_camera.h"
 #include "kernel_shader.h"
 #include "kernel_light.h"
 #include "kernel_emission.h"
 #include "kernel_random.h"
+#include "kernel_passes.h"
 
 CCL_NAMESPACE_BEGIN
 
-#ifdef __MODIFY_TP__
-__device float3 path_terminate_modified_throughput(KernelGlobals *kg, __global float3 *buffer, int x, int y, int offset, int stride, int sample)
-{
-	/* modify throughput to influence path termination probability, to avoid
-	   darker regions receiving fewer samples than lighter regions. also RGB
-	   are weighted differently. proper validation still remains to be done. */
-	const float3 weights = make_float3(1.0f, 1.33f, 0.66f);
-	const float3 one = make_float3(1.0f, 1.0f, 1.0f);
-	const int minsample = 5;
-	const float minL = 0.1f;
-
-	if(sample >= minsample) {
-		float3 L = buffer[offset + x + y*stride];
-		float3 Lmin = make_float3(minL, minL, minL);
-		float correct = (float)(sample+1)/(float)sample;
-
-		L = film_map(L*correct, sample);
-
-		return weights/clamp(L, Lmin, one);
-	}
-
-	return weights;
-}
-#endif
-
 typedef struct PathState {
 	uint flag;
 	int bounce;
@@ -168,7 +145,7 @@ __device_inline float path_state_terminate_probability(KernelGlobals *kg, PathSt
 	return average(throughput);
 }
 
-__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, float3 *light_L)
+__device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ray, Intersection *isect, BsdfEval *L_light)
 {
 	if(ray->t == 0.0f)
 		return false;
@@ -208,7 +185,7 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra
 				}
 
 				if(!scene_intersect(kg, ray, PATH_RAY_SHADOW_TRANSPARENT, isect)) {
-					*light_L *= throughput;
+					bsdf_eval_mul(L_light, throughput);
 					return false;
 				}
 
@@ -234,11 +211,14 @@ __device_inline bool shadow_blocked(KernelGlobals *kg, PathState *state, Ray *ra
 	return result;
 }
 
-__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, float3 throughput)
+__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, Ray ray, __global float *buffer)
 {
 	/* initialize */
-	float3 L = make_float3(0.0f, 0.0f, 0.0f);
-	float Ltransparent = 0.0f;
+	PathRadiance L;
+	float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
+	float L_transparent = 0.0f;
+
+	path_radiance_init(&L, kernel_data.film.use_light_pass);
 
 #ifdef __EMISSION__
 	float ray_pdf = 0.0f;
@@ -257,12 +237,12 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 		if(!scene_intersect(kg, &ray, visibility, &isect)) {
 			/* eval background shader if nothing hit */
 			if(kernel_data.background.transparent && (state.flag & PATH_RAY_CAMERA)) {
-				Ltransparent += average(throughput);
+				L_transparent += average(throughput);
 			}
 			else {
 				/* sample background shader */
-				float3 background_L = indirect_background(kg, &ray, state.flag, ray_pdf);
-				L += throughput*background_L;
+				float3 L_background = indirect_background(kg, &ray, state.flag, ray_pdf);
+				path_radiance_accum_background(&L, throughput, L_background, state.bounce);
 			}
 
 			break;
@@ -274,19 +254,24 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 		float rbsdf = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF);
 		shader_eval_surface(kg, &sd, rbsdf, state.flag);
 
+		kernel_write_data_passes(kg, buffer, &L, &sd, sample, state.flag, throughput);
+
 #ifdef __HOLDOUT__
 		if((sd.flag & SD_HOLDOUT) && (state.flag & PATH_RAY_CAMERA)) {
 			float3 holdout_weight = shader_holdout_eval(kg, &sd);
 
 			if(kernel_data.background.transparent)
-				Ltransparent += average(holdout_weight*throughput);
+				/* any throughput is ok, should all be identical here */
+				L_transparent += average(holdout_weight*throughput);
 		}
 #endif
 
 #ifdef __EMISSION__
 		/* emission */
-		if(sd.flag & SD_EMISSION)
-			L += throughput*indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf);
+		if(sd.flag & SD_EMISSION) {
+			float3 emission = indirect_emission(kg, &sd, isect.t, state.flag, ray_pdf);
+			path_radiance_accum_emission(&L, throughput, emission, state.bounce);
+		}
 #endif
 
 		/* path termination. this is a strange place to put the termination, it's
@@ -310,7 +295,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 				float light_v = path_rng(kg, rng, sample, rng_offset + PRNG_LIGHT_V);
 
 				Ray light_ray;
-				float3 light_L;
+				BsdfEval L_light;
 
 #ifdef __MULTI_LIGHT__
 				/* index -1 means randomly sample from distribution */
@@ -320,10 +305,10 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 #else
 				const int i = -1;
 #endif
-					if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &light_L)) {
+					if(direct_emission(kg, &sd, i, light_t, light_o, light_u, light_v, &light_ray, &L_light)) {
 						/* trace shadow ray */
-						if(!shadow_blocked(kg, &state, &light_ray, &isect, &light_L))
-							L += throughput*light_L;
+						if(!shadow_blocked(kg, &state, &light_ray, &isect, &L_light))
+							path_radiance_accum_light(&L, throughput, &L_light, state.bounce);
 					}
 #ifdef __MULTI_LIGHT__
 				}
@@ -338,7 +323,7 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 
 		/* sample BSDF */
 		float bsdf_pdf;
-		float3 bsdf_eval;
+		BsdfEval bsdf_eval;
 		float3 bsdf_omega_in;
 		differential3 bsdf_domega_in;
 		float bsdf_u = path_rng(kg, rng, sample, rng_offset + PRNG_BSDF_U);
@@ -350,11 +335,11 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 
 		shader_release(kg, &sd);
 
-		if(bsdf_pdf == 0.0f || is_zero(bsdf_eval))
+		if(bsdf_pdf == 0.0f || bsdf_eval_is_zero(&bsdf_eval))
 			break;
 
 		/* modify throughput */
-		throughput *= bsdf_eval/bsdf_pdf;
+		path_radiance_bsdf_bounce(&L, &throughput, &bsdf_eval, bsdf_pdf, state.bounce, label);
 
 		/* set labels */
 #if defined(__EMISSION__) || defined(__BACKGROUND__)
@@ -374,18 +359,33 @@ __device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int sample, R
 #endif
 	}
 
-	return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent);
+	float3 L_sum = path_radiance_sum(&L);
+
+	kernel_write_light_passes(kg, buffer, &L, sample);
+
+	return make_float4(L_sum.x, L_sum.y, L_sum.z, 1.0f - L_transparent);
 }
 
-__device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int sample, int x, int y, int offset, int stride)
+__device void kernel_path_trace(KernelGlobals *kg,
+	__global float *buffer, __global uint *rng_state,
+	int sample, int x, int y, int offset, int stride)
 {
+	/* buffer offset */
+	int index = offset + x + y*stride;
+	int pass_stride = kernel_data.film.pass_stride;
+
+	rng_state += index;
+	buffer += index*pass_stride;
+
+	kernel_clear_passes(buffer, sample, pass_stride);
+
 	/* initialize random numbers */
 	RNG rng;
 
 	float filter_u;
 	float filter_v;
 
-	path_rng_init(kg, rng_state, sample, &rng, x, y, offset, stride, &filter_u, &filter_v);
+	path_rng_init(kg, rng_state, sample, &rng, x, y, &filter_u, &filter_v);
 
 	/* sample camera ray */
 	Ray ray;
@@ -396,23 +396,12 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
 	camera_sample(kg, x, y, filter_u, filter_v, lens_u, lens_v, &ray);
 
 	/* integrate */
-#ifdef __MODIFY_TP__
-	float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, offset, stride, sample);
-	float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput)/throughput;
-#else
-	float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
-	float4 L = kernel_path_integrate(kg, &rng, sample, ray, throughput);
-#endif
+	float4 L = kernel_path_integrate(kg, &rng, sample, ray, buffer);
 
 	/* accumulate result in output buffer */
-	int index = offset + x + y*stride;
-
-	if(sample == 0)
-		buffer[index] = L;
-	else
-		buffer[index] += L;
+	kernel_write_pass_float4(buffer, sample, L);
 
-	path_rng_end(kg, rng_state, rng, x, y, offset, stride);
+	path_rng_end(kg, rng_state, rng);
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_random.h b/intern/cycles/kernel/kernel_random.h
index 41301ebd3dc..6d15100f8a3 100644
--- a/intern/cycles/kernel/kernel_random.h
+++ b/intern/cycles/kernel/kernel_random.h
@@ -123,7 +123,7 @@ __device_inline float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dime
 #endif
 }
 
-__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
+__device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
 {
 #ifdef __SOBOL_FULL_SCREEN__
 	uint px, py;
@@ -135,19 +135,31 @@ __device_inline void path_rng_init(KernelGlobals *kg, __global uint *rng_state,
 
 	*rng ^= kernel_data.integrator.seed;
 
-	*fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x;
-	*fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y;
+	if(sample == 0) {
+		*fx = 0.5f;
+		*fy = 0.5f;
+	}
+	else {
+		*fx = size * (float)px * (1.0f/(float)0xFFFFFFFF) - x;
+		*fy = size * (float)py * (1.0f/(float)0xFFFFFFFF) - y;
+	}
 #else
-	*rng = rng_state[offset + x + y*stride];
+	*rng = *rng_state;
 
 	*rng ^= kernel_data.integrator.seed;
 
-	*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
-	*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+	if(sample == 0) {
+		*fx = 0.5f;
+		*fy = 0.5f;
+	}
+	else {
+		*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
+		*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+	}
 #endif
 }
 
-__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
+__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng)
 {
 	/* nothing to do */
 }
@@ -163,21 +175,27 @@ __device float path_rng(KernelGlobals *kg, RNG *rng, int sample, int dimension)
 	return (float)*rng * (1.0f/(float)0xFFFFFFFF);
 }
 
-__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, int offset, int stride, float *fx, float *fy)
+__device void path_rng_init(KernelGlobals *kg, __global uint *rng_state, int sample, RNG *rng, int x, int y, float *fx, float *fy)
 {
 	/* load state */
-	*rng = rng_state[offset + x + y*stride];
+	*rng = *rng_state;
 
 	*rng ^= kernel_data.integrator.seed;
 
-	*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
-	*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+	if(sample == 0) {
+		*fx = 0.5f;
+		*fy = 0.5f;
+	}
+	else {
+		*fx = path_rng(kg, rng, sample, PRNG_FILTER_U);
+		*fy = path_rng(kg, rng, sample, PRNG_FILTER_V);
+	}
 }
 
-__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng, int x, int y, int offset, int stride)
+__device void path_rng_end(KernelGlobals *kg, __global uint *rng_state, RNG rng)
 {
 	/* store state for next sample */
-	rng_state[offset + x + y*stride] = rng;
+	*rng_state = rng;
 }
 
 #endif
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index 9c59e1566a9..1d2cf46aa56 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -75,7 +75,7 @@ __device_inline void shader_setup_from_ray(KernelGlobals *kg, ShaderData *sd,
 	if(sd->shader & SHADER_SMOOTH_NORMAL)
 		sd->N = triangle_smooth_normal(kg, sd->prim, sd->u, sd->v);
 
-	sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK);
+	sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2);
 
 #ifdef __DPDU__
 	/* dPdu/dPdv */
@@ -166,7 +166,7 @@ __device void shader_setup_from_sample(KernelGlobals *kg, ShaderData *sd,
 #endif
 	}
 
-	sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK);
+	sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2);
 
 #ifdef __DPDU__
 	/* dPdu/dPdv */
@@ -243,7 +243,7 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData
 	sd->Ng = -sd->P;
 	sd->I = -sd->P;
 	sd->shader = kernel_data.background.shader;
-	sd->flag = kernel_tex_fetch(__shader_flag, sd->shader & SHADER_MASK);
+	sd->flag = kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2);
 
 #ifdef __INSTANCING__
 	sd->object = ~0;
@@ -275,8 +275,8 @@ __device_inline void shader_setup_from_background(KernelGlobals *kg, ShaderData
 
 #ifdef __MULTI_CLOSURE__
 
-__device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf,
-	int skip_bsdf, float3 sum_eval, float sum_pdf, float sum_sample_weight)
+__device_inline void _shader_bsdf_multi_eval(const ShaderData *sd, const float3 omega_in, float *pdf,
+	int skip_bsdf, BsdfEval *bsdf_eval, float sum_pdf, float sum_sample_weight)
 {
 	for(int i = 0; i< sd->num_closure; i++) {
 		if(i == skip_bsdf)
@@ -293,7 +293,7 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float
 #endif
 
 			if(bsdf_pdf != 0.0f) {
-				sum_eval += eval*sc->weight;
+				bsdf_eval_accum(bsdf_eval, sc->type, eval*sc->weight);
 				sum_pdf += bsdf_pdf*sc->sample_weight;
 			}
 
@@ -302,25 +302,27 @@ __device_inline float3 _shader_bsdf_multi_eval(const ShaderData *sd, const float
 	}
 
 	*pdf = sum_pdf/sum_sample_weight;
-	return sum_eval;
 }
 
 #endif
 
-__device float3 shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd,
-	const float3 omega_in, float *pdf)
+__device void shader_bsdf_eval(KernelGlobals *kg, const ShaderData *sd,
+	const float3 omega_in, BsdfEval *eval, float *pdf)
 {
 #ifdef __MULTI_CLOSURE__
-	return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, make_float3(0.0f, 0.0f, 0.0f), 0.0f, 0.0f);
+	bsdf_eval_init(eval, NBUILTIN_CLOSURES, make_float3(0.0f, 0.0f, 0.0f), kernel_data.film.use_light_pass);
+
+	return _shader_bsdf_multi_eval(sd, omega_in, pdf, -1, eval, 0.0f, 0.0f);
 #else
 	const ShaderClosure *sc = &sd->closure;
+
 	*pdf = 0.0f;
-	return svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight;
+	*eval = svm_bsdf_eval(sd, sc, omega_in, pdf)*sc->weight;
 #endif
 }
 
 __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd,
-	float randu, float randv, float3 *eval,
+	float randu, float randv, BsdfEval *bsdf_eval,
 	float3 *omega_in, differential3 *domega_in, float *pdf)
 {
 #ifdef __MULTI_CLOSURE__
@@ -359,27 +361,28 @@ __device int shader_bsdf_sample(KernelGlobals *kg, const ShaderData *sd,
 
 	const ShaderClosure *sc = &sd->closure[sampled];
 	int label;
+	float3 eval;
 
 	*pdf = 0.0f;
 #ifdef __OSL__
-	label = OSLShader::bsdf_sample(sd, sc, randu, randv, *eval, *omega_in, *domega_in, *pdf);
+	label = OSLShader::bsdf_sample(sd, sc, randu, randv, eval, *omega_in, *domega_in, *pdf);
 #else
-	label = svm_bsdf_sample(sd, sc, randu, randv, eval, omega_in, domega_in, pdf);
+	label = svm_bsdf_sample(sd, sc, randu, randv, &eval, omega_in, domega_in, pdf);
 #endif
 
-	*eval *= sc->weight;
+	bsdf_eval_init(bsdf_eval, sc->type, eval*sc->weight, kernel_data.film.use_light_pass);
 
 	if(sd->num_closure > 1 && *pdf != 0.0f) {
 		float sweight = sc->sample_weight;
-		*eval = _shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, *eval, *pdf*sweight, sweight);
+		_shader_bsdf_multi_eval(sd, *omega_in, pdf, sampled, bsdf_eval, *pdf*sweight, sweight);
 	}
 
 	return label;
 #else
 	/* sample the single closure that we picked */
 	*pdf = 0.0f;
-	int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, eval, omega_in, domega_in, pdf);
-	*eval *= sd->closure.weight;
+	int label = svm_bsdf_sample(sd, &sd->closure, randu, randv, bsdf_eval, omega_in, domega_in, pdf);
+	*bsdf_eval *= sd->closure.weight;
 	return label;
 #endif
 }
@@ -421,6 +424,68 @@ __device float3 shader_bsdf_transparency(KernelGlobals *kg, ShaderData *sd)
 #endif
 }
 
+__device float3 shader_bsdf_diffuse(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef __MULTI_CLOSURE__
+	float3 eval = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i< sd->num_closure; i++) {
+		ShaderClosure *sc = &sd->closure[i];
+
+		if(CLOSURE_IS_BSDF_DIFFUSE(sc->type))
+			eval += sc->weight;
+	}
+
+	return eval;
+#else
+	if(CLOSURE_IS_BSDF_DIFFUSE(sd->closure.type))
+		return sd->closure.weight;
+	else
+		return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
+
+__device float3 shader_bsdf_glossy(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef __MULTI_CLOSURE__
+	float3 eval = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i< sd->num_closure; i++) {
+		ShaderClosure *sc = &sd->closure[i];
+
+		if(CLOSURE_IS_BSDF_GLOSSY(sc->type))
+			eval += sc->weight;
+	}
+
+	return eval;
+#else
+	if(CLOSURE_IS_BSDF_GLOSSY(sd->closure.type))
+		return sd->closure.weight;
+	else
+		return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
+
+__device float3 shader_bsdf_transmission(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef __MULTI_CLOSURE__
+	float3 eval = make_float3(0.0f, 0.0f, 0.0f);
+
+	for(int i = 0; i< sd->num_closure; i++) {
+		ShaderClosure *sc = &sd->closure[i];
+
+		if(CLOSURE_IS_BSDF_TRANSMISSION(sc->type))
+			eval += sc->weight;
+	}
+
+	return eval;
+#else
+	if(CLOSURE_IS_BSDF_TRANSMISSION(sd->closure.type))
+		return sd->closure.weight;
+	else
+		return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+}
 
 /* Emission */
 
@@ -588,12 +653,17 @@ __device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect)
 	int prim = kernel_tex_fetch(__prim_index, isect->prim);
 	float4 Ns = kernel_tex_fetch(__tri_normal, prim);
 	int shader = __float_as_int(Ns.w);
-	int flag = kernel_tex_fetch(__shader_flag, shader & SHADER_MASK);
+	int flag = kernel_tex_fetch(__shader_flag, (shader & SHADER_MASK)*2);
 
 	return (flag & SD_HAS_SURFACE_TRANSPARENT) != 0;
 }
 #endif
 
+__device int shader_pass_id(KernelGlobals *kg, ShaderData *sd)
+{
+	return kernel_tex_fetch(__shader_flag, (sd->shader & SHADER_MASK)*2 + 1);
+}
+
 /* Free ShaderData */
 
 __device void shader_release(KernelGlobals *kg, ShaderData *sd)
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 008ec0bdf28..b4b1da83162 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -70,6 +70,9 @@ CCL_NAMESPACE_BEGIN
 #ifdef __KERNEL_ADV_SHADING__
 #define __MULTI_CLOSURE__
 #define __TRANSPARENT_SHADOWS__
+#ifdef __KERNEL_CPU__
+#define __PASSES__
+#endif
 #endif
 
 //#define __MULTI_LIGHT__
@@ -150,6 +153,75 @@ typedef enum ClosureLabel {
 	LABEL_STOP = 2048
 } ClosureLabel;
 
+/* Render Passes */
+
+typedef enum PassType {
+	PASS_NONE = 0,
+	PASS_COMBINED = 1,
+	PASS_DEPTH = 2,
+	PASS_NORMAL = 8,
+	PASS_UV = 16,
+	PASS_OBJECT_ID = 32,
+	PASS_MATERIAL_ID = 64,
+	PASS_DIFFUSE_COLOR = 128,
+	PASS_GLOSSY_COLOR = 256,
+	PASS_TRANSMISSION_COLOR = 512,
+	PASS_DIFFUSE_INDIRECT = 1024,
+	PASS_GLOSSY_INDIRECT = 2048,
+	PASS_TRANSMISSION_INDIRECT = 4096,
+	PASS_DIFFUSE_DIRECT = 8192,
+	PASS_GLOSSY_DIRECT = 16384,
+	PASS_TRANSMISSION_DIRECT = 32768,
+	PASS_EMISSION = 65536,
+	PASS_BACKGROUND = 131072
+} PassType;
+
+#define PASS_ALL (~0)
+
+#ifdef __PASSES__
+
+typedef float3 PathThroughput;
+
+struct PathRadiance {
+	int use_light_pass;
+
+	float3 indirect;
+	float3 direct_throughput;
+	float3 direct_emission;
+
+	float3 color_diffuse;
+	float3 color_glossy;
+	float3 color_transmission;
+
+	float3 direct_diffuse;
+	float3 direct_glossy;
+	float3 direct_transmission;
+
+	float3 indirect_diffuse;
+	float3 indirect_glossy;
+	float3 indirect_transmission;
+
+	float3 emission;
+	float3 background;
+};
+
+struct BsdfEval {
+	int use_light_pass;
+
+	float3 diffuse;
+	float3 glossy;
+	float3 transmission;
+	float3 transparent;
+};
+
+#else
+
+typedef float3 PathThroughput;
+typedef float3 PathRadiance;
+typedef float3 BsdfEval;
+
+#endif
+
 /* Shader Flag */
 
 typedef enum ShaderFlag {
@@ -353,7 +425,32 @@ typedef struct KernelCamera {
 
 typedef struct KernelFilm {
 	float exposure;
-	int pad1, pad2, pad3;
+	int pass_flag;
+	int pass_stride;
+	int use_light_pass;
+
+	int pass_combined;
+	int pass_depth;
+	int pass_normal;
+	int pass_pad;
+
+	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_diffuse_indirect;
+	int pass_glossy_indirect;
+
+	int pass_transmission_indirect;
+	int pass_diffuse_direct;
+	int pass_glossy_direct;
+	int pass_transmission_direct;
+
+	int pass_emission;
+	int pass_background;
 } KernelFilm;
 
 typedef struct KernelBackground {
diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h
index ff6e9b94a0f..533a2944557 100644
--- a/intern/cycles/kernel/svm/svm_types.h
+++ b/intern/cycles/kernel/svm/svm_types.h
@@ -266,22 +266,26 @@ typedef enum ShaderType {
 
 typedef enum ClosureType {
 	CLOSURE_BSDF_ID,
+
 	CLOSURE_BSDF_DIFFUSE_ID,
 	CLOSURE_BSDF_OREN_NAYAR_ID,
-	CLOSURE_BSDF_TRANSLUCENT_ID,
+
 	CLOSURE_BSDF_REFLECTION_ID,
-	CLOSURE_BSDF_REFRACTION_ID,
-	CLOSURE_BSDF_GLASS_ID,
-	CLOSURE_BSDF_TRANSPARENT_ID,
 	CLOSURE_BSDF_MICROFACET_GGX_ID,
-	CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID,
 	CLOSURE_BSDF_MICROFACET_BECKMANN_ID,
-	CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID,
 	CLOSURE_BSDF_WARD_ID,
-	CLOSURE_BSDF_ASHIKHMIN_VELVET_ID,
-	CLOSURE_BSDF_WESTIN_BACKSCATTER_ID,
 	CLOSURE_BSDF_WESTIN_SHEEN_ID,
 
+	CLOSURE_BSDF_TRANSLUCENT_ID,
+	CLOSURE_BSDF_REFRACTION_ID,
+	CLOSURE_BSDF_WESTIN_BACKSCATTER_ID,
+	CLOSURE_BSDF_ASHIKHMIN_VELVET_ID,
+	CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID,
+	CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID,
+	CLOSURE_BSDF_GLASS_ID,
+
+	CLOSURE_BSDF_TRANSPARENT_ID,
+
 	CLOSURE_BSSRDF_CUBIC_ID,
 	CLOSURE_EMISSION_ID,
 	CLOSURE_DEBUG_ID,
@@ -297,7 +301,10 @@ typedef enum ClosureType {
 } ClosureType;
 
 /* watch this, being lazy with memory usage */
-#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_WESTIN_SHEEN_ID)
+#define CLOSURE_IS_BSDF(type) (type <= CLOSURE_BSDF_TRANSPARENT_ID)
+#define CLOSURE_IS_BSDF_DIFFUSE(type) (type >= CLOSURE_BSDF_DIFFUSE_ID && type <= CLOSURE_BSDF_OREN_NAYAR_ID)
+#define CLOSURE_IS_BSDF_GLOSSY(type) (type >= CLOSURE_BSDF_REFLECTION_ID && type <= CLOSURE_BSDF_WESTIN_SHEEN_ID)
+#define CLOSURE_IS_BSDF_TRANSMISSION(type) (type >= CLOSURE_BSDF_TRANSLUCENT_ID && type <= CLOSURE_BSDF_GLASS_ID)
 #define CLOSURE_IS_VOLUME(type) (type >= CLOSURE_VOLUME_ID && type <= CLOSURE_VOLUME_ISOTROPIC_ID)
 #define CLOSURE_IS_EMISSION(type) (type == CLOSURE_EMISSION_ID)
 #define CLOSURE_IS_HOLDOUT(type) (type == CLOSURE_HOLDOUT_ID)