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

29 files changed, 1225 insertions, 187 deletions

diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 763cff0bd7d..6353bf37a15 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -183,6 +183,38 @@ class CyclesRender_PT_layers(CyclesButtonsPanel, Panel):
 
         layout.separator()
 
+        split = layout.split()
+
+        col = split.column()
+        col.label(text="Passes:")
+        col.prop(rl, "use_pass_combined")
+        col.prop(rl, "use_pass_z")
+        col.prop(rl, "use_pass_normal")
+        col.prop(rl, "use_pass_object_index")
+        col.prop(rl, "use_pass_material_index")
+        col.prop(rl, "use_pass_emit")
+        col.prop(rl, "use_pass_environment")
+
+        col = split.column()
+        col.label()
+        col.label(text="Diffuse:")
+        row = col.row(align=True)
+        row.prop(rl, "use_pass_diffuse_direct", text="Direct", toggle=True)
+        row.prop(rl, "use_pass_diffuse_indirect", text="Indirect", toggle=True)
+        row.prop(rl, "use_pass_diffuse_color", text="Color", toggle=True)
+        col.label(text="Glossy:")
+        row = col.row(align=True)
+        row.prop(rl, "use_pass_glossy_direct", text="Direct", toggle=True)
+        row.prop(rl, "use_pass_glossy_indirect", text="Indirect", toggle=True)
+        row.prop(rl, "use_pass_glossy_color", text="Color", toggle=True)
+        col.label(text="Transmission:")
+        row = col.row(align=True)
+        row.prop(rl, "use_pass_transmission_direct", text="Direct", toggle=True)
+        row.prop(rl, "use_pass_transmission_indirect", text="Indirect", toggle=True)
+        row.prop(rl, "use_pass_transmission_color", text="Color", toggle=True)
+
+        layout.separator()
+
         rl = rd.layers[0]
         layout.prop(rl, "material_override", text="Material")
 
diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp
index c805bd03063..afcfcd95063 100644
--- a/intern/cycles/blender/blender_object.cpp
+++ b/intern/cycles/blender/blender_object.cpp
@@ -214,6 +214,7 @@ void BlenderSync::sync_object(BL::Object b_parent, int b_index, BL::Object b_ob,
 	/* object sync */
 	if(object_updated || (object->mesh && object->mesh->need_update)) {
 		object->name = b_ob.name().c_str();
+		object->pass_id = b_ob.pass_index();
 		object->tfm = tfm;
 
 		/* visibility flags for both parent */
diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp
index b052fb3e195..ff1c32831bb 100644
--- a/intern/cycles/blender/blender_session.cpp
+++ b/intern/cycles/blender/blender_session.cpp
@@ -116,9 +116,68 @@ void BlenderSession::free_session()
 	delete session;
 }
 
+static PassType get_pass_type(BL::RenderPass b_pass)
+{
+	switch(b_pass.type()) {
+		case BL::RenderPass::type_COMBINED:
+			return PASS_COMBINED;
+
+		case BL::RenderPass::type_Z:
+			return PASS_DEPTH;
+		case BL::RenderPass::type_NORMAL:
+			return PASS_NORMAL;
+		case BL::RenderPass::type_OBJECT_INDEX:
+			return PASS_OBJECT_ID;
+		case BL::RenderPass::type_UV:
+			return PASS_UV;
+		case BL::RenderPass::type_MATERIAL_INDEX:
+			return PASS_MATERIAL_ID;
+
+		case BL::RenderPass::type_DIFFUSE_DIRECT:
+			return PASS_DIFFUSE_DIRECT;
+		case BL::RenderPass::type_GLOSSY_DIRECT:
+			return PASS_GLOSSY_DIRECT;
+		case BL::RenderPass::type_TRANSMISSION_DIRECT:
+			return PASS_TRANSMISSION_DIRECT;
+
+		case BL::RenderPass::type_DIFFUSE_INDIRECT:
+			return PASS_DIFFUSE_INDIRECT;
+		case BL::RenderPass::type_GLOSSY_INDIRECT:
+			return PASS_GLOSSY_INDIRECT;
+		case BL::RenderPass::type_TRANSMISSION_INDIRECT:
+			return PASS_TRANSMISSION_INDIRECT;
+
+		case BL::RenderPass::type_DIFFUSE_COLOR:
+			return PASS_DIFFUSE_COLOR;
+		case BL::RenderPass::type_GLOSSY_COLOR:
+			return PASS_GLOSSY_COLOR;
+		case BL::RenderPass::type_TRANSMISSION_COLOR:
+			return PASS_TRANSMISSION_COLOR;
+
+		case BL::RenderPass::type_EMIT:
+			return PASS_EMISSION;
+		case BL::RenderPass::type_ENVIRONMENT:
+			return PASS_BACKGROUND;
+
+		case BL::RenderPass::type_DIFFUSE:
+		case BL::RenderPass::type_SHADOW:
+		case BL::RenderPass::type_AO:
+		case BL::RenderPass::type_COLOR:
+		case BL::RenderPass::type_REFRACTION:
+		case BL::RenderPass::type_SPECULAR:
+		case BL::RenderPass::type_REFLECTION:
+		case BL::RenderPass::type_VECTOR:
+		case BL::RenderPass::type_MIST:
+			return PASS_NONE;
+	}
+	
+	return PASS_NONE;
+}
+
 void BlenderSession::render()
 {
 	/* get buffer parameters */
+	SessionParams session_params = BlenderSync::get_session_params(b_userpref, b_scene, background);
 	BufferParams buffer_params = BlenderSync::get_buffer_params(b_scene, b_rv3d, width, height);
 	int w = buffer_params.width, h = buffer_params.height;
 
@@ -143,6 +202,25 @@ void BlenderSession::render()
 		/* set layer */
 		b_rlay = *b_iter;
 
+		/* add passes */
+		if(session_params.device.type == DEVICE_CPU) { /* todo */
+			BL::RenderLayer::passes_iterator b_pass_iter;
+			
+			for(b_rlay.passes.begin(b_pass_iter); b_pass_iter != b_rlay.passes.end(); ++b_pass_iter) {
+				BL::RenderPass b_pass(*b_pass_iter);
+				PassType pass_type = get_pass_type(b_pass);
+
+				if(pass_type != PASS_NONE)
+					Pass::add(pass_type, buffer_params.passes);
+			}
+		}
+
+		scene->film->passes = buffer_params.passes;
+		scene->film->need_update = true;
+
+		/* update session */
+		session->reset(buffer_params, session_params.samples);
+
 		/* update scene */
 		sync->sync_data(b_v3d, active);
 
@@ -165,22 +243,41 @@ void BlenderSession::write_render_result()
 {
 	/* get state */
 	RenderBuffers *buffers = session->buffers;
+
+	/* copy data from device */
+	if(!buffers->copy_from_device())
+		return;
+
+	BufferParams& params = buffers->params;
 	float exposure = scene->film->exposure;
 	double total_time, sample_time;
 	int sample;
+
 	session->progress.get_sample(sample, total_time, sample_time);
 
-	/* get pixels */
-	float4 *pixels = buffers->copy_from_device(exposure, sample);
+	vector<float> pixels(params.width*params.height*4);
 
-	if(!pixels)
-		return;
+	/* copy each pass */
+	BL::RenderLayer::passes_iterator b_iter;
+	
+	for(b_rlay.passes.begin(b_iter); b_iter != b_rlay.passes.end(); ++b_iter) {
+		BL::RenderPass b_pass(*b_iter);
 
-	/* write pixels */
-	rna_RenderLayer_rect_set(&b_rlay.ptr, (float*)pixels);
-	RE_engine_update_result((RenderEngine*)b_engine.ptr.data, (RenderResult*)b_rr.ptr.data);
+		/* find matching pass type */
+		PassType pass_type = get_pass_type(b_pass);
+		int components = b_pass.channels();
+
+		/* copy pixels */
+		if(buffers->get_pass(pass_type, exposure, sample, components, &pixels[0]))
+			rna_RenderPass_rect_set(&b_pass.ptr, &pixels[0]);
+	}
 
-	delete [] pixels;
+	/* copy combined pass */
+	if(buffers->get_pass(PASS_COMBINED, exposure, sample, 4, &pixels[0]))
+		rna_RenderLayer_rect_set(&b_rlay.ptr, &pixels[0]);
+
+	/* tag result as updated */
+	RE_engine_update_result((RenderEngine*)b_engine.ptr.data, (RenderResult*)b_rr.ptr.data);
 }
 
 void BlenderSession::synchronize()
diff --git a/intern/cycles/blender/blender_shader.cpp b/intern/cycles/blender/blender_shader.cpp
index f7b750e6dab..5310e35dc25 100644
--- a/intern/cycles/blender/blender_shader.cpp
+++ b/intern/cycles/blender/blender_shader.cpp
@@ -636,6 +636,7 @@ void BlenderSync::sync_materials()
 			ShaderGraph *graph = new ShaderGraph();
 
 			shader->name = b_mat->name().c_str();
+			shader->pass_id = b_mat->pass_index();
 
 			/* create nodes */
 			if(b_mat->use_nodes() && b_mat->node_tree()) {
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp
index 25da978edd8..2ca599f6c67 100644
--- a/intern/cycles/device/device_cpu.cpp
+++ b/intern/cycles/device/device_cpu.cpp
@@ -162,7 +162,7 @@ public:
 		if(system_cpu_support_optimized()) {
 			for(int y = task.y; y < task.y + task.h; y++) {
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_optimized_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
+					kernel_cpu_optimized_path_trace(kg, (float*)task.buffer, (unsigned int*)task.rng_state,
 						task.sample, x, y, task.offset, task.stride);
 
 				if(tasks.worker_cancel())
@@ -174,7 +174,7 @@ public:
 		{
 			for(int y = task.y; y < task.y + task.h; y++) {
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_path_trace(kg, (float4*)task.buffer, (unsigned int*)task.rng_state,
+					kernel_cpu_path_trace(kg, (float*)task.buffer, (unsigned int*)task.rng_state,
 						task.sample, x, y, task.offset, task.stride);
 
 				if(tasks.worker_cancel())
@@ -194,7 +194,7 @@ public:
 		if(system_cpu_support_optimized()) {
 			for(int y = task.y; y < task.y + task.h; y++)
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
+					kernel_cpu_optimized_tonemap(kg, (uchar4*)task.rgba, (float*)task.buffer,
 						task.sample, task.resolution, x, y, task.offset, task.stride);
 		}
 		else
@@ -202,7 +202,7 @@ public:
 		{
 			for(int y = task.y; y < task.y + task.h; y++)
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float4*)task.buffer,
+					kernel_cpu_tonemap(kg, (uchar4*)task.rgba, (float*)task.buffer,
 						task.sample, task.resolution, x, y, task.offset, task.stride);
 		}
 	}
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)
diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp
index a6bbbc91901..08dda944111 100644
--- a/intern/cycles/render/buffers.cpp
+++ b/intern/cycles/render/buffers.cpp
@@ -22,6 +22,7 @@
 #include "device.h"
 
 #include "util_debug.h"
+#include "util_foreach.h"
 #include "util_hash.h"
 #include "util_image.h"
 #include "util_math.h"
@@ -31,6 +32,48 @@
 
 CCL_NAMESPACE_BEGIN
 
+/* Buffer Params */
+
+BufferParams::BufferParams()
+{
+	width = 0;
+	height = 0;
+
+	full_x = 0;
+	full_y = 0;
+	full_width = 0;
+	full_height = 0;
+
+	Pass::add(PASS_COMBINED, passes);
+}
+
+void BufferParams::get_offset_stride(int& offset, int& stride)
+{
+	offset = -(full_x + full_y*width);
+	stride = width;
+}
+
+bool BufferParams::modified(const BufferParams& params)
+{
+	return !(full_x == params.full_x
+		&& full_y == params.full_y
+		&& width == params.width
+		&& height == params.height
+		&& full_width == params.full_width
+		&& full_height == params.full_height
+		&& Pass::equals(passes, params.passes));
+}
+
+int BufferParams::get_passes_size()
+{
+	int size = 0;
+
+	foreach(Pass& pass, passes)
+		size += pass.components;
+	
+	return size;
+}
+
 /* Render Buffers */
 
 RenderBuffers::RenderBuffers(Device *device_)
@@ -64,7 +107,7 @@ void RenderBuffers::reset(Device *device, BufferParams& params_)
 	device_free();
 	
 	/* allocate buffer */
-	buffer.resize(params.width, params.height);
+	buffer.resize(params.width*params.height*params.get_passes_size());
 	device->mem_alloc(buffer, MEM_READ_WRITE);
 	device->mem_zero(buffer);
 
@@ -82,31 +125,76 @@ void RenderBuffers::reset(Device *device, BufferParams& params_)
 	device->mem_copy_to(rng_state);
 }
 
-float4 *RenderBuffers::copy_from_device(float exposure, int sample)
+bool RenderBuffers::copy_from_device()
 {
 	if(!buffer.device_pointer)
-		return NULL;
+		return false;
 
 	device->mem_copy_from(buffer, 0, params.width, params.height, sizeof(float4));
 
-	float4 *out = new float4[params.width*params.height];
-	float4 *in = (float4*)buffer.data_pointer;
-	float scale = 1.0f/(float)sample;
-	
-	for(int i = params.width*params.height - 1; i >= 0; i--) {
-		float4 rgba = in[i]*scale;
+	return true;
+}
+
+bool RenderBuffers::get_pass(PassType type, float exposure, int sample, int components, float *pixels)
+{
+	int pass_offset = 0;
+
+	foreach(Pass& pass, params.passes) {
+		if(pass.type != type) {
+			pass_offset += pass.components;
+			continue;
+		}
+
+		float *in = (float*)buffer.data_pointer + pass_offset;
+		int pass_stride = params.get_passes_size();
+
+		float scale = (pass.filter)? 1.0f/(float)sample: 1.0f;
+		float scale_exposure = (pass.exposure)? scale*exposure: scale;
+
+		int size = params.width*params.height;
+
+		if(components == 1) {
+			assert(pass.components == components);
+
+			/* scalar */
+			for(int i = 0; i < size; i++, in += pass_stride, pixels++) {
+				float f = *in;
+
+				pixels[0] = f*scale_exposure;
+			}
+		}
+		else if(components == 3) {
+			assert(pass.components == 4);
+
+			/* RGB/vector */
+			for(int i = 0; i < size; i++, in += pass_stride, pixels += 3) {
+				float3 f = make_float3(in[0], in[1], in[2]);
+
+				pixels[0] = f.x*scale_exposure;
+				pixels[1] = f.y*scale_exposure;
+				pixels[2] = f.z*scale_exposure;
+			}
+		}
+		else if(components == 4) {
+			assert(pass.components == components);
+
+			/* RGBA */
+			for(int i = 0; i < size; i++, in += pass_stride, pixels += 4) {
+				float4 f = make_float4(in[0], in[1], in[2], in[3]);
 
-		rgba.x = rgba.x*exposure;
-		rgba.y = rgba.y*exposure;
-		rgba.z = rgba.z*exposure;
+				pixels[0] = f.x*scale_exposure;
+				pixels[1] = f.y*scale_exposure;
+				pixels[2] = f.z*scale_exposure;
 
-		/* clamp since alpha might be > 1.0 due to russian roulette */
-		rgba.w = clamp(rgba.w, 0.0f, 1.0f);
+				/* clamp since alpha might be > 1.0 due to russian roulette */
+				pixels[3] = clamp(f.w*scale, 0.0f, 1.0f);
+			}
+		}
 
-		out[i] = rgba;
+		return true;
 	}
 
-	return out;
+	return false;
 }
 
 /* Display Buffer */
diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h
index f4a9b37c09b..3062e5ae3e4 100644
--- a/intern/cycles/render/buffers.h
+++ b/intern/cycles/render/buffers.h
@@ -21,6 +21,10 @@
 
 #include "device_memory.h"
 
+#include "film.h"
+
+#include "kernel_types.h"
+
 #include "util_string.h"
 #include "util_thread.h"
 #include "util_types.h"
@@ -45,32 +49,16 @@ public:
 	int full_width;
 	int full_height;
 
-	BufferParams()
-	{
-		width = 0;
-		height = 0;
-
-		full_x = 0;
-		full_y = 0;
-		full_width = 0;
-		full_height = 0;
-	}
-
-	void get_offset_stride(int& offset, int& stride)
-	{
-		offset = -(full_x + full_y*width);
-		stride = width;
-	}
-
-	bool modified(const BufferParams& params)
-	{
-		return !(full_x == params.full_x
-			&& full_y == params.full_y
-			&& width == params.width
-			&& height == params.height
-			&& full_width == params.full_width
-			&& full_height == params.full_height);
-	}
+	/* passes */
+	vector<Pass> passes;
+
+	/* functions */
+	BufferParams();
+
+	void get_offset_stride(int& offset, int& stride);
+	bool modified(const BufferParams& params);
+	void add_pass(PassType type);
+	int get_passes_size();
 };
 
 /* Render Buffers */
@@ -80,7 +68,7 @@ public:
 	/* buffer parameters */
 	BufferParams params;
 	/* float buffer */
-	device_vector<float4> buffer;
+	device_vector<float> buffer;
 	/* random number generator state */
 	device_vector<uint> rng_state;
 	/* mutex, must be locked manually by callers */
@@ -90,7 +78,9 @@ public:
 	~RenderBuffers();
 
 	void reset(Device *device, BufferParams& params);
-	float4 *copy_from_device(float exposure, int sample);
+
+	bool copy_from_device();
+	bool get_pass(PassType type, float exposure, int sample, int components, float *pixels);
 
 protected:
 	void device_free();
diff --git a/intern/cycles/render/film.cpp b/intern/cycles/render/film.cpp
index 0ae2866f182..bc51384b873 100644
--- a/intern/cycles/render/film.cpp
+++ b/intern/cycles/render/film.cpp
@@ -21,11 +21,111 @@
 #include "film.h"
 #include "scene.h"
 
+#include "util_foreach.h"
+
 CCL_NAMESPACE_BEGIN
 
+/* Pass */
+
+void Pass::add(PassType type, vector<Pass>& passes)
+{
+	Pass pass;
+
+	pass.type = type;
+	pass.filter = true;
+	pass.exposure = false;
+
+	switch(type) {
+		case PASS_NONE:
+			pass.components = 0;
+			break;
+		case PASS_COMBINED:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_DEPTH:
+			pass.components = 1;
+			pass.filter = false;
+			break;
+		case PASS_NORMAL:
+			pass.components = 4;
+			break;
+		case PASS_UV:
+			pass.components = 4;
+			break;
+		case PASS_OBJECT_ID:
+			pass.components = 1;
+			pass.filter = false;
+			break;
+		case PASS_MATERIAL_ID:
+			pass.components = 1;
+			pass.filter = false;
+			break;
+		case PASS_DIFFUSE_COLOR:
+			pass.components = 4;
+			break;
+		case PASS_GLOSSY_COLOR:
+			pass.components = 4;
+			break;
+		case PASS_TRANSMISSION_COLOR:
+			pass.components = 4;
+			break;
+		case PASS_DIFFUSE_INDIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_GLOSSY_INDIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_TRANSMISSION_INDIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_DIFFUSE_DIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_GLOSSY_DIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_TRANSMISSION_DIRECT:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+
+		case PASS_EMISSION:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+		case PASS_BACKGROUND:
+			pass.components = 4;
+			pass.exposure = true;
+			break;
+	}
+
+	passes.push_back(pass);
+}
+
+bool Pass::equals(const vector<Pass>& A, const vector<Pass>& B)
+{
+	if(A.size() != B.size())
+		return false;
+	
+	for(int i = 0; i < A.size(); i++)
+		if(A[i].type != B[i].type)
+			return false;
+	
+	return true;
+}
+
+/* Film */
+
 Film::Film()
 {
 	exposure = 0.8f;
+	Pass::add(PASS_COMBINED, passes);
 	need_update = true;
 }
 
@@ -42,6 +142,82 @@ void Film::device_update(Device *device, DeviceScene *dscene)
 
 	/* update __data */
 	kfilm->exposure = exposure;
+	kfilm->pass_flag = 0;
+	kfilm->pass_stride = 0;
+	kfilm->use_light_pass = 0;
+
+	foreach(Pass& pass, passes) {
+		kfilm->pass_flag |= pass.type;
+
+		switch(pass.type) {
+			case PASS_COMBINED:
+				kfilm->pass_combined = kfilm->pass_stride;
+				break;
+			case PASS_DEPTH:
+				kfilm->pass_depth = kfilm->pass_stride;
+				break;
+			case PASS_NORMAL:
+				kfilm->pass_normal = kfilm->pass_stride;
+				break;
+			case PASS_UV:
+				kfilm->pass_uv = kfilm->pass_stride;
+				break;
+			case PASS_OBJECT_ID:
+				kfilm->pass_object_id = kfilm->pass_stride;
+				break;
+			case PASS_MATERIAL_ID:
+				kfilm->pass_material_id = kfilm->pass_stride;
+				break;
+			case PASS_DIFFUSE_COLOR:
+				kfilm->pass_diffuse_color = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_GLOSSY_COLOR:
+				kfilm->pass_glossy_color = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_TRANSMISSION_COLOR:
+				kfilm->pass_transmission_color = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_DIFFUSE_INDIRECT:
+				kfilm->pass_diffuse_indirect = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_GLOSSY_INDIRECT:
+				kfilm->pass_glossy_indirect = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_TRANSMISSION_INDIRECT:
+				kfilm->pass_transmission_indirect = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_DIFFUSE_DIRECT:
+				kfilm->pass_diffuse_direct = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_GLOSSY_DIRECT:
+				kfilm->pass_glossy_direct = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_TRANSMISSION_DIRECT:
+				kfilm->pass_transmission_direct = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+
+			case PASS_EMISSION:
+				kfilm->pass_emission = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+				break;
+			case PASS_BACKGROUND:
+				kfilm->pass_background = kfilm->pass_stride;
+				kfilm->use_light_pass = 1;
+			case PASS_NONE:
+				break;
+		}
+
+		kfilm->pass_stride += pass.components;
+	}
 
 	need_update = false;
 }
@@ -52,7 +228,8 @@ void Film::device_free(Device *device, DeviceScene *dscene)
 
 bool Film::modified(const Film& film)
 {
-	return !(exposure == film.exposure);
+	return !(exposure == film.exposure
+		&& Pass::equals(passes, film.passes));
 }
 
 void Film::tag_update(Scene *scene)
diff --git a/intern/cycles/render/film.h b/intern/cycles/render/film.h
index df24fad3725..511ad316460 100644
--- a/intern/cycles/render/film.h
+++ b/intern/cycles/render/film.h
@@ -20,6 +20,9 @@
 #define __FILM_H__
 
 #include "util_string.h"
+#include "util_vector.h"
+
+#include "kernel_types.h"
 
 CCL_NAMESPACE_BEGIN
 
@@ -27,9 +30,21 @@ class Device;
 class DeviceScene;
 class Scene;
 
+class Pass {
+public:
+	PassType type;
+	int components;
+	bool filter;
+	bool exposure;
+
+	static void add(PassType type, vector<Pass>& passes);
+	static bool equals(const vector<Pass>& A, const vector<Pass>& B);
+};
+
 class Film {
 public:
 	float exposure;
+	vector<Pass> passes;
 	bool need_update;
 
 	Film();
diff --git a/intern/cycles/render/object.cpp b/intern/cycles/render/object.cpp
index 3a9f0add735..f83c85c632d 100644
--- a/intern/cycles/render/object.cpp
+++ b/intern/cycles/render/object.cpp
@@ -36,6 +36,7 @@ Object::Object()
 	mesh = NULL;
 	tfm = transform_identity();
 	visibility = ~0;
+	pass_id = 0;
 }
 
 Object::~Object()
@@ -135,6 +136,7 @@ void ObjectManager::device_update_transforms(Device *device, DeviceScene *dscene
 		/* todo: correct for displacement, and move to a better place */
 		float uniform_scale;
 		float surface_area = 0.0f;
+		float pass_id = ob->pass_id;
 		
 		if(transform_uniform_scale(tfm, uniform_scale)) {
 			map<Mesh*, float>::iterator it = surface_area_map.find(mesh);
@@ -171,7 +173,7 @@ void ObjectManager::device_update_transforms(Device *device, DeviceScene *dscene
 		memcpy(&objects[offset], &tfm, sizeof(float4)*4);
 		memcpy(&objects[offset+4], &itfm, sizeof(float4)*4);
 		memcpy(&objects[offset+8], &ntfm, sizeof(float4)*4);
-		objects[offset+12] = make_float4(surface_area, 0.0f, 0.0f, 0.0f);
+		objects[offset+12] = make_float4(surface_area, pass_id, 0.0f, 0.0f);
 
 		i++;
 
diff --git a/intern/cycles/render/object.h b/intern/cycles/render/object.h
index 7fe83cf7d91..14da2cfb35d 100644
--- a/intern/cycles/render/object.h
+++ b/intern/cycles/render/object.h
@@ -41,6 +41,7 @@ public:
 	Transform tfm;
 	BoundBox bounds;
 	ustring name;
+	int pass_id;
 	vector<ParamValue> attributes;
 	uint visibility;
 
diff --git a/intern/cycles/render/shader.cpp b/intern/cycles/render/shader.cpp
index 6e827ec94bb..12968a79ab2 100644
--- a/intern/cycles/render/shader.cpp
+++ b/intern/cycles/render/shader.cpp
@@ -35,6 +35,7 @@ CCL_NAMESPACE_BEGIN
 Shader::Shader()
 {
 	name = "";
+	pass_id = 0;
 
 	graph = NULL;
 	graph_bump = NULL;
@@ -167,7 +168,7 @@ void ShaderManager::device_update_common(Device *device, DeviceScene *dscene, Sc
 	if(scene->shaders.size() == 0)
 		return;
 
-	uint shader_flag_size = scene->shaders.size()*2;
+	uint shader_flag_size = scene->shaders.size()*4;
 	uint *shader_flag = dscene->shader_flag.resize(shader_flag_size);
 	uint i = 0;
 
@@ -184,7 +185,9 @@ void ShaderManager::device_update_common(Device *device, DeviceScene *dscene, Sc
 			flag |= SD_HOMOGENEOUS_VOLUME;
 
 		shader_flag[i++] = flag;
+		shader_flag[i++] = shader->pass_id;
 		shader_flag[i++] = flag;
+		shader_flag[i++] = shader->pass_id;
 	}
 
 	device->tex_alloc("__shader_flag", dscene->shader_flag);
diff --git a/intern/cycles/render/shader.h b/intern/cycles/render/shader.h
index 45efa123ef6..35f3cfe27f5 100644
--- a/intern/cycles/render/shader.h
+++ b/intern/cycles/render/shader.h
@@ -47,6 +47,7 @@ class Shader {
 public:
 	/* name */
 	string name;
+	int pass_id;
 
 	/* shader graph */
 	ShaderGraph *graph;