From bae896691aa3d7bb2a75292da3cc490894996b01 Mon Sep 17 00:00:00 2001
From: Brecht Van Lommel <brechtvanlommel@pandora.be>
Date: Sun, 28 Aug 2011 13:55:59 +0000
Subject: Cycles: * Add alpha pass output, to use set Transparent option in
 Film panel. * Add Holdout closure (OSL terminology), this is like the Sky
 option in the   internal renderer, objects with this closure show the
 background / zero   alpha. * Add option to use Gaussian instead of Box pixel
 filter in the UI. * Remove camera response curves for now, they don't really
 belong here in   the pipeline, should be moved to compositor.

* Output full float values for rendering now, previously was only byte precision.
* Add a patch from Thomas to get a preview passes option, but still disabled
  because it isn't quite working right yet.
* CUDA: don't compile shader graph evaluation inline.
* Convert tabs to spaces in python files.
---
 intern/cycles/app/cycles_test.cpp                  |   8 +-
 intern/cycles/blender/addon/__init__.py            |  74 +--
 intern/cycles/blender/addon/engine.py              |  62 +-
 intern/cycles/blender/addon/enums.py               |  81 +--
 intern/cycles/blender/addon/properties.py          | 186 +++---
 intern/cycles/blender/addon/ui.py                  | 733 +++++++++++----------
 intern/cycles/blender/addon/xml.py                 |  90 +--
 intern/cycles/blender/blender_session.cpp          |  48 +-
 intern/cycles/blender/blender_shader.cpp           |   7 +
 intern/cycles/blender/blender_sync.cpp             |  26 +-
 intern/cycles/cmake/external_libs.cmake            |   8 +-
 intern/cycles/device/device.cpp                    |  10 +-
 intern/cycles/device/device.h                      |   2 +-
 intern/cycles/device/device_cuda.cpp               |  14 +-
 intern/cycles/device/device_multi.cpp              |   4 +-
 intern/cycles/kernel/kernel.cpp                    |  12 -
 intern/cycles/kernel/kernel_compat_cuda.h          |   1 +
 intern/cycles/kernel/kernel_compat_opencl.h        |   1 +
 intern/cycles/kernel/kernel_film.h                 |  24 +-
 intern/cycles/kernel/kernel_path.h                 |  65 +-
 intern/cycles/kernel/kernel_shader.h               |  20 +
 intern/cycles/kernel/kernel_textures.h             |   3 -
 intern/cycles/kernel/kernel_types.h                |  15 +-
 intern/cycles/kernel/osl/osl_shader.cpp            |  19 +-
 intern/cycles/kernel/osl/osl_shader.h              |   1 +
 intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h   |   2 +-
 intern/cycles/kernel/svm/bsdf_diffuse.h            |   4 +-
 intern/cycles/kernel/svm/bsdf_microfacet.h         |   4 +-
 intern/cycles/kernel/svm/bsdf_reflection.h         |   1 +
 intern/cycles/kernel/svm/bsdf_refraction.h         |   1 +
 intern/cycles/kernel/svm/bsdf_transparent.h        |   1 +
 intern/cycles/kernel/svm/bsdf_ward.h               |   2 +-
 intern/cycles/kernel/svm/bsdf_westin.h             |   4 +-
 intern/cycles/kernel/svm/svm.h                     |   4 +-
 intern/cycles/kernel/svm/svm_closure.h             |   6 +
 intern/cycles/kernel/svm/svm_convert.h             |   2 +-
 intern/cycles/kernel/svm/svm_types.h               |   3 +-
 intern/cycles/render/CMakeLists.txt                |   2 -
 intern/cycles/render/background.cpp                |   6 +-
 intern/cycles/render/background.h                  |   1 +
 intern/cycles/render/buffers.cpp                   |  59 +-
 intern/cycles/render/buffers.h                     |   4 +
 intern/cycles/render/camera.h                      |   2 -
 intern/cycles/render/film.cpp                      |  72 +-
 intern/cycles/render/film.h                        |   2 -
 intern/cycles/render/filter.cpp                    |  11 +
 intern/cycles/render/filter.h                      |   7 +
 intern/cycles/render/mesh.cpp                      |   2 +-
 intern/cycles/render/nodes.cpp                     |  18 +
 intern/cycles/render/nodes.h                       |   5 +
 intern/cycles/render/scene.h                       |   5 -
 intern/cycles/render/session.cpp                   |  49 +-
 intern/cycles/render/session.h                     |  10 +-
 intern/cycles/render/tile.cpp                      |  12 +-
 intern/cycles/render/tile.h                        |   3 +-
 intern/cycles/util/util_color.h                    |   5 +
 intern/cycles/util/util_types.h                    |   1 +
 source/blender/blenkernel/BKE_node.h               |   1 +
 source/blender/blenkernel/intern/node.c            |   3 +-
 .../blender/makesrna/intern/rna_nodetree_types.h   |   3 +-
 source/blender/nodes/CMakeLists.txt                |   1 +
 source/blender/nodes/SHD_node.h                    |   1 +
 .../blender/nodes/intern/SHD_nodes/SHD_holdout.c   |  63 ++
 63 files changed, 1037 insertions(+), 859 deletions(-)
 create mode 100644 source/blender/nodes/intern/SHD_nodes/SHD_holdout.c

diff --git a/intern/cycles/app/cycles_test.cpp b/intern/cycles/app/cycles_test.cpp
index f617e5db44b..c16f6f4e99a 100644
--- a/intern/cycles/app/cycles_test.cpp
+++ b/intern/cycles/app/cycles_test.cpp
@@ -85,7 +85,7 @@ static void session_print_status()
 static void session_init()
 {
 	options.session = new Session(options.session_params);
-	options.session->reset(options.width, options.height);
+	options.session->reset(options.width, options.height, options.session_params.passes);
 	options.session->scene = options.scene;
 	
 	if(options.session_params.background && !options.quiet)
@@ -162,13 +162,13 @@ static void resize(int width, int height)
 	options.height= height;
 
 	if(options.session)
-		options.session->reset(options.width, options.height);
+		options.session->reset(options.width, options.height, options.session_params.passes);
 }
 
 void keyboard(unsigned char key)
 {
 	if(key == 'r')
-		options.session->reset(options.width, options.height);
+		options.session->reset(options.width, options.height, options.session_params.passes);
 	else if(key == 27) // escape
 		options.session->progress.set_cancel("Cancelled");
 }
@@ -285,7 +285,7 @@ using namespace ccl;
 
 int main(int argc, const char **argv)
 {
-	path_init("../blender/intern/cycles");
+	path_init("../build/bin/2.59/scripts/addons/cycles/");
 
 	options_parse(argc, argv);
 
diff --git a/intern/cycles/blender/addon/__init__.py b/intern/cycles/blender/addon/__init__.py
index e66d078f8c7..418f3d4cd27 100644
--- a/intern/cycles/blender/addon/__init__.py
+++ b/intern/cycles/blender/addon/__init__.py
@@ -37,50 +37,50 @@ from cycles import xml
 from cycles import engine
 
 class CyclesRender(bpy.types.RenderEngine):
-	bl_idname = 'CYCLES'
-	bl_label = "Cycles"
+    bl_idname = 'CYCLES'
+    bl_label = "Cycles"
 
-	def __init__(self):
-		engine.init()
-		self.session = None
-	
-	def __del__(self):
-		engine.free(self)
+    def __init__(self):
+        engine.init()
+        self.session = None
+    
+    def __del__(self):
+        engine.free(self)
 
-	# final render
-	def update(self, data, scene):
-		engine.create(self, data, scene)
-		engine.update(self, data, scene)
+    # final render
+    def update(self, data, scene):
+        engine.create(self, data, scene)
+        engine.update(self, data, scene)
 
-	def render(self):
-		engine.render(self)
+    def render(self):
+        engine.render(self)
 
-	# preview render
-	# def preview_update(self, context, id):
-	#	pass
-	#
-	# def preview_render(self):
-	#	pass
-	
-	# viewport render
-	def view_update(self, context):
-		if not self.session:
-			engine.create(self, context.blend_data, context.scene,
-				context.region, context.space_data, context.region_data)
-		engine.update(self, context.blend_data, context.scene)
+    # preview render
+    # def preview_update(self, context, id):
+    #    pass
+    #
+    # def preview_render(self):
+    #    pass
+    
+    # viewport render
+    def view_update(self, context):
+        if not self.session:
+            engine.create(self, context.blend_data, context.scene,
+                context.region, context.space_data, context.region_data)
+        engine.update(self, context.blend_data, context.scene)
 
-	def view_draw(self, context):
-		engine.draw(self, context.region, context.space_data, context.region_data)
+    def view_draw(self, context):
+        engine.draw(self, context.region, context.space_data, context.region_data)
 
 def register():
-	properties.register()
-	ui.register()
-	xml.register()
-	bpy.utils.register_module(__name__)
+    properties.register()
+    ui.register()
+    xml.register()
+    bpy.utils.register_module(__name__)
 
 def unregister():
-	xml.unregister()
-	ui.unregister()
-	properties.unregister()
-	bpy.utils.unregister_module(__name__)
+    xml.unregister()
+    ui.unregister()
+    properties.unregister()
+    bpy.utils.unregister_module(__name__)
 
diff --git a/intern/cycles/blender/addon/engine.py b/intern/cycles/blender/addon/engine.py
index d25eb21eeb9..d6ea15a435f 100644
--- a/intern/cycles/blender/addon/engine.py
+++ b/intern/cycles/blender/addon/engine.py
@@ -19,52 +19,52 @@
 import bpy
 
 def init():
-	import libcycles_blender as lib
-	import os.path
-	lib.init(os.path.dirname(__file__))
+    import libcycles_blender as lib
+    import os.path
+    lib.init(os.path.dirname(__file__))
 
 def create(engine, data, scene, region = 0, v3d = 0, rv3d = 0):
-	import libcycles_blender as lib
+    import libcycles_blender as lib
 
-	data = data.as_pointer()
-	scene = scene.as_pointer()
-	if region:
-		region = region.as_pointer()
-	if v3d:
-		v3d = v3d.as_pointer()
-	if rv3d:
-		rv3d = rv3d.as_pointer()
+    data = data.as_pointer()
+    scene = scene.as_pointer()
+    if region:
+        region = region.as_pointer()
+    if v3d:
+        v3d = v3d.as_pointer()
+    if rv3d:
+        rv3d = rv3d.as_pointer()
 
-	engine.session = lib.create(engine.as_pointer(), data, scene, region, v3d, rv3d)
+    engine.session = lib.create(engine.as_pointer(), data, scene, region, v3d, rv3d)
 
 def free(engine):
-	if "session" in dir(engine):
-		if engine.session:
-			import libcycles_blender as lib
-			lib.free(engine.session)
-		del engine.session
+    if "session" in dir(engine):
+        if engine.session:
+            import libcycles_blender as lib
+            lib.free(engine.session)
+        del engine.session
 
 def render(engine):
-	import libcycles_blender as lib
-	lib.render(engine.session)
+    import libcycles_blender as lib
+    lib.render(engine.session)
 
 def update(engine, data, scene):
-	import libcycles_blender as lib
-	lib.sync(engine.session)
+    import libcycles_blender as lib
+    lib.sync(engine.session)
 
 def draw(engine, region, v3d, rv3d):
-	import libcycles_blender as lib
-	v3d = v3d.as_pointer()
-	rv3d = rv3d.as_pointer()
+    import libcycles_blender as lib
+    v3d = v3d.as_pointer()
+    rv3d = rv3d.as_pointer()
 
-	# draw render image
-	lib.draw(engine.session, v3d, rv3d)
+    # draw render image
+    lib.draw(engine.session, v3d, rv3d)
 
 def available_devices():
-	import libcycles_blender as lib
-	return lib.available_devices()
+    import libcycles_blender as lib
+    return lib.available_devices()
 
 def with_osl():
-	import libcycles_blender as lib
-	return lib.with_osl()
+    import libcycles_blender as lib
+    return lib.with_osl()
 
diff --git a/intern/cycles/blender/addon/enums.py b/intern/cycles/blender/addon/enums.py
index fd12fa6d5a5..4b30f606de2 100644
--- a/intern/cycles/blender/addon/enums.py
+++ b/intern/cycles/blender/addon/enums.py
@@ -33,81 +33,8 @@ bvh_types = (
 ("DYNAMIC_BVH", "Dynamic BVH", "Objects can be individually updated, at the cost of slower render time"),
 ("STATIC_BVH", "Static BVH", "Any object modification requires a complete BVH rebuild, but renders faster"))
 
-response_curves = (
-("None", "None", ""),
-("", "Agfa", ""),
-("Agfacolor Futura 100", "Futura 100", ""),
-("Agfacolor Futura 200", "Futura 200", ""),
-("Agfacolor Futura 400", "Futura 400", ""),
-("Agfacolor Futura II 100", "Futura II 100", ""),
-("Agfacolor Futura II 200", "Futura II 200", ""),
-("Agfacolor Futura II 400", "Futura II 400", ""),
-("Agfacolor HDC 100 plus", "HDC 100 plus", ""),
-("Agfacolor HDC 400 plus", "HDC 400 plus", ""),
-("Agfacolor HDC 200 plus", "HDC 200 plus", ""),
-("Agfacolor Optima II 100", "Optima II 100", ""),
-("Agfacolor Optima II 200", "Optima II 200", ""),
-("Agfacolor Ultra 050", "Ultra 050", ""),
-("", "Agfa", ""),
-("Agfacolor Vista 100", "Vista 100", ""),
-("Agfacolor Vista 200", "Vista 200", ""),
-("Agfacolor Vista 400", "Vista 400", ""),
-("Agfacolor Vista 800", "Vista 800", ""),
-("Agfachrome CT Precisa 100", "CT Precisa 100", ""),
-("Agfachrome CT Precisa 200", "CT Precisa 200", ""),
-("Agfachrome RSX2 050", "Agfachrome RSX2 050", ""),
-("Agfachrome RSX2 100", "Agfachrome RSX2 100", ""),
-("Agfachrome RSX2 200", "Agfachrome RSX2 200", ""),
-("Advantix 100", "Advantix 100", ""),
-("Advantix 200", "Advantix 200", ""),
-("Advantix 400", "Advantix 400", ""),
-("", "Kodak", ""),
-("Gold 100", "Gold 100", ""),
-("Gold 200", "Gold 200", ""),
-("Max Zoom 800", "Max Zoom 800", ""),
-("Portra 100T", "Portra 100T", ""),
-("Portra 160NC", "Portra 160NC", ""),
-("Portra 160VC", "Portra 160VC", ""),
-("Portra 800", "Portra 800", ""),
-("Portra 400VC", "Portra 400VC", ""),
-("Portra 400NC", "Portra 400NC", ""),
-("", "Kodak", ""),
-("Ektachrome 100 plus", "Ektachrome 100 plus", ""),
-("Ektachrome 320T", "Ektachrome 320T", ""),
-("Ektachrome 400X", "Ektachrome 400X", ""),
-("Ektachrome 64", "Ektachrome 64", ""),
-("Ektachrome 64T", "Ektachrome 64T", ""),
-("Ektachrome E100S", "Ektachrome E100S", ""),
-("Ektachrome 100", "Ektachrome 100", ""),
-("Kodachrome 200", "Kodachrome 200", ""),
-("Kodachrome 25", "Kodachrome 25", ""),
-("Kodachrome 64", "Kodachrome 64", ""),
-#("DSCS 3151", "DSCS 3151", ""),
-#("DSCS 3152", "DSCS 3152", ""),
-#("DSCS 3153", "DSCS 3153", ""),
-#("DSCS 3154", "DSCS 3154", ""),
-#("DSCS 3155", "DSCS 3155", ""),
-#("DSCS 3156", "DSCS 3156", ""),
-#("KAI-0311", "KAI-0311", ""),
-#("KAF-2001", "KAF-2001", ""),
-#("KAF-3000", "KAF-3000", ""),
-#("KAI-0372", "KAI-0372", ""),
-#("KAI-1010", "KAI-1010", ""),
-("", "Fujifilm", ""),
-("F-125", "F-125", ""),
-("F-250", "F-250", ""),
-("F-400", "F-400", ""),
-("FCI", "FCI", ""),
-("FP2900Z", "FP2900Z", ""),
-("", "Eastman", ""),
-("Double X Neg 12min", "Double X Neg 12min", ""),
-("Double X Neg 6min", "Double X Neg 6min", ""),
-("Double X Neg 5min", "Double X Neg 5min", ""),
-("Double X Neg 4min", "Double X Neg 4min", ""),
-("", "Canon", ""),
-("Optura 981111", "Optura 981111", ""),
-("Optura 981113", "Optura 981113", ""),
-("Optura 981114", "Optura 981114", ""),
-("Optura 981111.SLRR", "Optura 981111.SLRR", "")
-)
+filter_types = (
+("BOX", "Box", "Box filter"),
+("GAUSSIAN", "Gaussian", "Gaussian filter"))
+
 
diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index 3ec0587d89c..62c1db5a16d 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -22,99 +22,117 @@ from bpy.props import *
 from cycles import enums
 
 class CyclesRenderSettings(bpy.types.PropertyGroup):
-	@classmethod
-	def register(cls):
-		bpy.types.Scene.cycles = PointerProperty(type=cls, name="Cycles Render Settings", description="Cycles Render Settings")
-
-		cls.device = EnumProperty(name="Device", description="Device to use for rendering",
-			items=enums.devices, default="CPU")
-
-		cls.shading_system = EnumProperty(name="Shading System", description="Shading system to use for rendering",
-			items=enums.shading_systems, default="GPU_COMPATIBLE")
-
-		cls.passes = IntProperty(name="Passes", description="Number of passes to render",
-			default=10, min=1, max=2147483647)
-		cls.min_bounces = IntProperty(name="Min Bounces", description="Minimum number of bounces",
-			default=3, min=0, max=1024)
-		cls.max_bounces = IntProperty(name="Max Bounces", description="Maximum number of bounces",
-			default=8, min=0, max=1024)
-		cls.no_caustics = BoolProperty(name="No Caustics", description="Leave out caustics, resulting in a darker image with less noise",
-			default=False)
-		cls.blur_caustics = FloatProperty(name="Blur Caustics", description="Blur caustics to reduce noise",
-			default=0.0, min=0.0, max=1.0)
-
-		cls.exposure = FloatProperty(name="Exposure", description="Image brightness scale",
-			default=1.0, min=0.0, max=10.0)
-		cls.response_curve = EnumProperty(name="Response Curve", description="Measured camera film response",
-			items=enums.response_curves, default="Advantix 400")
-
-		cls.debug_tile_size = IntProperty(name="Tile Size", description="",
-			default=1024, min=1, max=4096)
-		cls.debug_min_size = IntProperty(name="Min Size", description="",
-			default=64, min=1, max=4096)
-		cls.debug_reset_timeout = FloatProperty(name="Reset timeout", description="",
-			default=0.1, min=0.01, max=10.0)
-		cls.debug_cancel_timeout = FloatProperty(name="Cancel timeout", description="",
-			default=0.1, min=0.01, max=10.0)
-		cls.debug_text_timeout = FloatProperty(name="Text timeout", description="",
-			default=1.0, min=0.01, max=10.0)
-
-		cls.debug_bvh_type = EnumProperty(name="BVH Type", description="Choose between faster updates, or faster render",
-			items=enums.bvh_types, default="DYNAMIC_BVH")
-		cls.debug_use_spatial_splits = BoolProperty(name="Use Spatial Splits", description="Use BVH spatial splits: longer builder time, faster render",
-			default=False)
-
-	@classmethod
-	def unregister(cls):
-		del bpy.types.Scene.cycles
+    @classmethod
+    def register(cls):
+        bpy.types.Scene.cycles = PointerProperty(type=cls, name="Cycles Render Settings", description="Cycles Render Settings")
+
+        cls.device = EnumProperty(name="Device", description="Device to use for rendering",
+            items=enums.devices, default="CPU")
+
+        cls.shading_system = EnumProperty(name="Shading System", description="Shading system to use for rendering",
+            items=enums.shading_systems, default="GPU_COMPATIBLE")
+
+        cls.passes = IntProperty(name="Passes", description="Number of passes to render",
+            default=10, min=1, max=2147483647)
+        cls.preview_passes = IntProperty(name="Preview Passes", description="Number of passes to render in the viewport, unlimited if 0",
+            default=0, min=0, max=2147483647)
+        cls.min_bounces = IntProperty(name="Min Bounces", description="Minimum number of bounces",
+            default=3, min=0, max=1024)
+        cls.max_bounces = IntProperty(name="Max Bounces", description="Maximum number of bounces",
+            default=8, min=0, max=1024)
+        cls.no_caustics = BoolProperty(name="No Caustics", description="Leave out caustics, resulting in a darker image with less noise",
+            default=False)
+        cls.blur_caustics = FloatProperty(name="Blur Caustics", description="Blur caustics to reduce noise",
+            default=0.0, min=0.0, max=1.0)
+
+        cls.exposure = FloatProperty(name="Exposure", description="Image brightness scale",
+            default=1.0, min=0.0, max=10.0)
+        cls.transparent = BoolProperty(name="Transparent", description="World background is transparent",
+            default=False)
+
+        cls.filter_type = EnumProperty(name="Filter Type", description="Pixel filter type",
+            items=enums.filter_types, default="GAUSSIAN")
+        cls.filter_width = FloatProperty(name="Filter Width", description="Pixel filter width",
+            default=1.5, min=0.01, max=10.0)
+
+        cls.debug_tile_size = IntProperty(name="Tile Size", description="",
+            default=1024, min=1, max=4096)
+        cls.debug_min_size = IntProperty(name="Min Size", description="",
+            default=64, min=1, max=4096)
+        cls.debug_reset_timeout = FloatProperty(name="Reset timeout", description="",
+            default=0.1, min=0.01, max=10.0)
+        cls.debug_cancel_timeout = FloatProperty(name="Cancel timeout", description="",
+            default=0.1, min=0.01, max=10.0)
+        cls.debug_text_timeout = FloatProperty(name="Text timeout", description="",
+            default=1.0, min=0.01, max=10.0)
+
+        cls.debug_bvh_type = EnumProperty(name="BVH Type", description="Choose between faster updates, or faster render",
+            items=enums.bvh_types, default="DYNAMIC_BVH")
+        cls.debug_use_spatial_splits = BoolProperty(name="Use Spatial Splits", description="Use BVH spatial splits: longer builder time, faster render",
+            default=False)
+
+    @classmethod
+    def unregister(cls):
+        del bpy.types.Scene.cycles
 
 class CyclesCameraSettings(bpy.types.PropertyGroup):
-	@classmethod
-	def register(cls):
-		bpy.types.Camera.cycles = PointerProperty(type=cls, name="Cycles Camera Settings", description="Cycles Camera Settings")
+    @classmethod
+    def register(cls):
+        bpy.types.Camera.cycles = PointerProperty(type=cls, name="Cycles Camera Settings", description="Cycles Camera Settings")
 
-		cls.lens_radius = FloatProperty(name="Lens radius", description="Lens radius for depth of field",
-			default=0.0, min=0.0, max=10.0)
-	
-	@classmethod
-	def unregister(cls):
-		del bpy.types.Camera.cycles
+        cls.lens_radius = FloatProperty(name="Lens radius", description="Lens radius for depth of field",
+            default=0.0, min=0.0, max=10.0)
+    
+    @classmethod
+    def unregister(cls):
+        del bpy.types.Camera.cycles
 
 class CyclesMaterialSettings(bpy.types.PropertyGroup):
-	@classmethod
-	def register(cls):
-		bpy.types.Material.cycles = PointerProperty(type=cls, name="Cycles Material Settings", description="Cycles Material Settings")
+    @classmethod
+    def register(cls):
+        bpy.types.Material.cycles = PointerProperty(type=cls, name="Cycles Material Settings", description="Cycles Material Settings")
 
-	@classmethod
-	def unregister(cls):
-		del bpy.types.Material.cycles
+    @classmethod
+    def unregister(cls):
+        del bpy.types.Material.cycles
+
+class CyclesWorldSettings(bpy.types.PropertyGroup):
+    @classmethod
+    def register(cls):
+        bpy.types.World.cycles = PointerProperty(type=cls, name="Cycles World Settings", description="Cycles World Settings")
+
+    @classmethod
+    def unregister(cls):
+        del bpy.types.World.cycles
 
 class CyclesMeshSettings(bpy.types.PropertyGroup):
-	@classmethod
-	def register(cls):
-		bpy.types.Mesh.cycles = PointerProperty(type=cls, name="Cycles Mesh Settings", description="Cycles Mesh Settings")
-		bpy.types.Curve.cycles = PointerProperty(type=cls, name="Cycles Mesh Settings", description="Cycles Mesh Settings")
-		bpy.types.MetaBall.cycles = PointerProperty(type=cls, name="Cycles Mesh Settings", description="Cycles Mesh Settings")
-
-		cls.displacement_method = EnumProperty(name="Displacement Method", description="Method to use for the displacement",
-			items=enums.displacement_methods, default="BUMP")
-		cls.use_subdivision = BoolProperty(name="Use Subdivision", description="Subdivide mesh for rendering",
-			default=False)
-		cls.dicing_rate = FloatProperty(name="Dicing Rate", description="", default=1.0, min=0.001, max=1000.0)
-
-	@classmethod
-	def unregister(cls):
-		del bpy.types.Mesh.cycles
+    @classmethod
+    def register(cls):
+        bpy.types.Mesh.cycles = PointerProperty(type=cls, name="Cycles Mesh Settings", description="Cycles Mesh Settings")
+        bpy.types.Curve.cycles = PointerProperty(type=cls, name="Cycles Mesh Settings", description="Cycles Mesh Settings")
+        bpy.types.MetaBall.cycles = PointerProperty(type=cls, name="Cycles Mesh Settings", description="Cycles Mesh Settings")
+
+        cls.displacement_method = EnumProperty(name="Displacement Method", description="Method to use for the displacement",
+            items=enums.displacement_methods, default="BUMP")
+        cls.use_subdivision = BoolProperty(name="Use Subdivision", description="Subdivide mesh for rendering",
+            default=False)
+        cls.dicing_rate = FloatProperty(name="Dicing Rate", description="", default=1.0, min=0.001, max=1000.0)
+
+    @classmethod
+    def unregister(cls):
+        del bpy.types.Mesh.cycles
 
 def register():
-	bpy.utils.register_class(CyclesRenderSettings)
-	bpy.utils.register_class(CyclesCameraSettings)
-	bpy.utils.register_class(CyclesMaterialSettings)
-	bpy.utils.register_class(CyclesMeshSettings)
-	
+    bpy.utils.register_class(CyclesRenderSettings)
+    bpy.utils.register_class(CyclesCameraSettings)
+    bpy.utils.register_class(CyclesMaterialSettings)
+    bpy.utils.register_class(CyclesWorldSettings)
+    bpy.utils.register_class(CyclesMeshSettings)
+    
 def unregister():
-	bpy.utils.unregister_class(CyclesRenderSettings)
-	bpy.utils.unregister_class(CyclesCameraSettings)
-	bpy.utils.unregister_class(CyclesMaterialSettings)
-	bpy.utils.unregister_class(CyclesMeshSettings)
+    bpy.utils.unregister_class(CyclesRenderSettings)
+    bpy.utils.unregister_class(CyclesCameraSettings)
+    bpy.utils.unregister_class(CyclesMaterialSettings)
+    bpy.utils.unregister_class(CyclesWorldSettings)
+    bpy.utils.unregister_class(CyclesMeshSettings)
 
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index ca839691111..a1363ece854 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -22,494 +22,503 @@ from cycles import enums
 from cycles import engine
 
 class CyclesButtonsPanel():
-	bl_space_type = "PROPERTIES"
-	bl_region_type = "WINDOW"
-	bl_context = "render"
-	
-	@classmethod
-	def poll(cls, context):
-		rd = context.scene.render
-		return rd.engine == 'CYCLES'
+    bl_space_type = "PROPERTIES"
+    bl_region_type = "WINDOW"
+    bl_context = "render"
+    
+    @classmethod
+    def poll(cls, context):
+        rd = context.scene.render
+        return rd.engine == 'CYCLES'
 
 class CyclesRender_PT_integrator(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Integrator"
+    bl_label = "Integrator"
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		scene = context.scene
-		cycles = scene.cycles
+        scene = context.scene
+        cscene = scene.cycles
 
-		split = layout.split()
+        split = layout.split()
 
-		col = split.column()
-		col.prop(cycles, "passes")
-		col.prop(cycles, "no_caustics")
+        col = split.column()
+        col.prop(cscene, "passes", text="Render Passes")
+        #sub = col.row()
+        #sub.active = cscene.preview_passes >= 1
+        #sub.prop(cscene, "preview_passes")
+        col.prop(cscene, "no_caustics")
 
-		col = split.column()
-		col = col.column(align=True)
-		col.prop(cycles, "max_bounces")
-		col.prop(cycles, "min_bounces")
+        col = split.column()
+        col = col.column(align=True)
+        col.prop(cscene, "max_bounces")
+        col.prop(cscene, "min_bounces")
 
-		#row = col.row()
-		#row.prop(cycles, "blur_caustics")
-		#row.active = not cycles.no_caustics
-		
+        #row = col.row()
+        #row.prop(cscene, "blur_caustics")
+        #row.active = not cscene.no_caustics
+        
 class CyclesRender_PT_film(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Film"
+    bl_label = "Film"
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		scene = context.scene
-		cycles = scene.cycles
+        scene = context.scene
+        cscene = scene.cycles
 
-		split = layout.split()
+        split = layout.split()
 
-		split.prop(cycles, "exposure")
-		split.prop(cycles, "response_curve", text="")
+        col = split.column();
+        col.prop(cscene, "exposure")
+        col.prop(cscene, "transparent")
+
+        col = split.column()
+        col.prop(cscene, "filter_type", text="")
+        if cscene.filter_type != 'BOX':
+            col.prop(cscene, "filter_width", text="Width")
 
 class CyclesRender_PT_performance(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Performance"
-	bl_options = {'DEFAULT_CLOSED'}
+    bl_label = "Performance"
+    bl_options = {'DEFAULT_CLOSED'}
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		scene = context.scene
-		rd = scene.render
-		cycles = scene.cycles
+        scene = context.scene
+        rd = scene.render
+        cscene = scene.cycles
 
-		split = layout.split()
+        split = layout.split()
 
-		col = split.column(align=True)
+        col = split.column(align=True)
 
-		col.label(text="Threads:")
-		col.row().prop(rd, "threads_mode", expand=True)
-		sub = col.column()
-		sub.enabled = rd.threads_mode == 'FIXED'
-		sub.prop(rd, "threads")
+        col.label(text="Threads:")
+        col.row().prop(rd, "threads_mode", expand=True)
+        sub = col.column()
+        sub.enabled = rd.threads_mode == 'FIXED'
+        sub.prop(rd, "threads")
 
-		sub = col.column(align=True)
-		sub.label(text="Tiles:")
-		sub.prop(cycles, "debug_tile_size")
-		sub.prop(cycles, "debug_min_size")
+        sub = col.column(align=True)
+        sub.label(text="Tiles:")
+        sub.prop(cscene, "debug_tile_size")
+        sub.prop(cscene, "debug_min_size")
 
-		col = split.column()
+        col = split.column()
 
-		sub = col.column(align=True)
-		sub.label(text="Acceleration structure:")
-		sub.prop(cycles, "debug_bvh_type", text="")
-		sub.prop(cycles, "debug_use_spatial_splits")
+        sub = col.column(align=True)
+        sub.label(text="Acceleration structure:")
+        sub.prop(cscene, "debug_bvh_type", text="")
+        sub.prop(cscene, "debug_use_spatial_splits")
 
 class Cycles_PT_post_processing(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Post Processing"
-	bl_options = {'DEFAULT_CLOSED'}
+    bl_label = "Post Processing"
+    bl_options = {'DEFAULT_CLOSED'}
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		rd = context.scene.render
+        rd = context.scene.render
 
-		split = layout.split()
+        split = layout.split()
 
-		col = split.column()
-		col.prop(rd, "use_compositing")
-		col.prop(rd, "use_sequencer")
+        col = split.column()
+        col.prop(rd, "use_compositing")
+        col.prop(rd, "use_sequencer")
 
-		col = split.column()
-		col.prop(rd, "dither_intensity", text="Dither", slider=True)
+        col = split.column()
+        col.prop(rd, "dither_intensity", text="Dither", slider=True)
 
 class Cycles_PT_camera(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Cycles"
-	bl_context = "data"
+    bl_label = "Cycles"
+    bl_context = "data"
 
-	@classmethod
-	def poll(cls, context):
-		return context.camera
+    @classmethod
+    def poll(cls, context):
+        return context.camera
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		camera = context.camera
-		cycles = camera.cycles
+        camera = context.camera
+        ccamera = camera.cycles
 
-		layout.prop(cycles, "lens_radius")
+        layout.prop(ccamera, "lens_radius")
 
 class Cycles_PT_context_material(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Surface"
-	bl_context = "material"
-	bl_options = {'HIDE_HEADER'}
+    bl_label = "Surface"
+    bl_context = "material"
+    bl_options = {'HIDE_HEADER'}
 
-	@classmethod
-	def poll(cls, context):
-		return (context.material or context.object) and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        return (context.material or context.object) and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		mat = context.material
-		ob = context.object
-		slot = context.material_slot
-		space = context.space_data
+        mat = context.material
+        ob = context.object
+        slot = context.material_slot
+        space = context.space_data
 
-		if ob:
-			row = layout.row()
+        if ob:
+            row = layout.row()
 
-			row.template_list(ob, "material_slots", ob, "active_material_index", rows=2)
+            row.template_list(ob, "material_slots", ob, "active_material_index", rows=2)
 
-			col = row.column(align=True)
-			col.operator("object.material_slot_add", icon='ZOOMIN', text="")
-			col.operator("object.material_slot_remove", icon='ZOOMOUT', text="")
+            col = row.column(align=True)
+            col.operator("object.material_slot_add", icon='ZOOMIN', text="")
+            col.operator("object.material_slot_remove", icon='ZOOMOUT', text="")
 
-			col.menu("MATERIAL_MT_specials", icon='DOWNARROW_HLT', text="")
+            col.menu("MATERIAL_MT_specials", icon='DOWNARROW_HLT', text="")
 
-			if ob.mode == 'EDIT':
-				row = layout.row(align=True)
-				row.operator("object.material_slot_assign", text="Assign")
-				row.operator("object.material_slot_select", text="Select")
-				row.operator("object.material_slot_deselect", text="Deselect")
+            if ob.mode == 'EDIT':
+                row = layout.row(align=True)
+                row.operator("object.material_slot_assign", text="Assign")
+                row.operator("object.material_slot_select", text="Select")
+                row.operator("object.material_slot_deselect", text="Deselect")
 
-		split = layout.split(percentage=0.65)
+        split = layout.split(percentage=0.65)
 
-		if ob:
-			split.template_ID(ob, "active_material", new="material.new")
-			row = split.row()
+        if ob:
+            split.template_ID(ob, "active_material", new="material.new")
+            row = split.row()
 
-			if slot:
-				row.prop(slot, "link", text="")
-			else:
-				row.label()
-		elif mat:
-			split.template_ID(space, "pin_id")
-			split.separator()
+            if slot:
+                row.prop(slot, "link", text="")
+            else:
+                row.label()
+        elif mat:
+            split.template_ID(space, "pin_id")
+            split.separator()
 
 class Cycles_PT_mesh_displacement(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Displacement"
-	bl_context = "data"
+    bl_label = "Displacement"
+    bl_context = "data"
 
-	@classmethod
-	def poll(cls, context):
-		return context.mesh or context.curve or context.meta_ball
+    @classmethod
+    def poll(cls, context):
+        return context.mesh or context.curve or context.meta_ball
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		mesh = context.mesh
-		curve = context.curve
-		mball = context.meta_ball
+        mesh = context.mesh
+        curve = context.curve
+        mball = context.meta_ball
 
-		if mesh:
-			cycles = mesh.cycles
-		elif curve:
-			cycles = curve.cycles
-		elif mball:
-			cycles = mball.cycles
+        if mesh:
+            cdata = mesh.cycles
+        elif curve:
+            cdata = curve.cycles
+        elif mball:
+            cdata = mball.cycles
 
-		layout.prop(cycles, "displacement_method", text="Method")
-		layout.prop(cycles, "use_subdivision");
-		layout.prop(cycles, "dicing_rate");
+        layout.prop(cdata, "displacement_method", text="Method")
+        layout.prop(cdata, "use_subdivision");
+        layout.prop(cdata, "dicing_rate");
 
 def find_node(material, nodetype):
-	if material and material.node_tree:
-		ntree = material.node_tree
+    if material and material.node_tree:
+        ntree = material.node_tree
 
-		for node in ntree.nodes:
-			if type(node) is not bpy.types.NodeGroup and node.type == nodetype:
-				return node
-	
-	return None
+        for node in ntree.nodes:
+            if type(node) is not bpy.types.NodeGroup and node.type == nodetype:
+                return node
+    
+    return None
 
 def find_node_input(node, name):
-	for input in node.inputs:
-		if input.name == name:
-			return input
-	
-	return None
+    for input in node.inputs:
+        if input.name == name:
+            return input
+    
+    return None
 
 def panel_node_draw(layout, id, output_type, input_name):
-	if not id.node_tree:
-		layout.prop(id, "use_nodes")
-		return
+    if not id.node_tree:
+        layout.prop(id, "use_nodes")
+        return
 
-	ntree = id.node_tree
+    ntree = id.node_tree
 
-	node = find_node(id, output_type)
-	if not node:
-		layout.label(text="No output node.")
-	else:
-		input = find_node_input(node, input_name)
-		layout.template_node_view(ntree, node, input);
+    node = find_node(id, output_type)
+    if not node:
+        layout.label(text="No output node.")
+    else:
+        input = find_node_input(node, input_name)
+        layout.template_node_view(ntree, node, input);
 
 class CyclesLamp_PT_lamp(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Surface"
-	bl_context = "data"
+    bl_label = "Surface"
+    bl_context = "data"
 
-	@classmethod
-	def poll(cls, context):
-		return context.lamp and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        return context.lamp and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		mat = context.lamp
-		panel_node_draw(layout, mat, 'OUTPUT_LAMP', 'Surface')
+        mat = context.lamp
+        panel_node_draw(layout, mat, 'OUTPUT_LAMP', 'Surface')
 
 class CyclesWorld_PT_surface(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Surface"
-	bl_context = "world"
+    bl_label = "Surface"
+    bl_context = "world"
 
-	@classmethod
-	def poll(cls, context):
-		return context.world and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        return context.world and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		mat = context.world
-		panel_node_draw(layout, mat, 'OUTPUT_WORLD', 'Surface')
+        mat = context.world
+        panel_node_draw(layout, mat, 'OUTPUT_WORLD', 'Surface')
 
 class CyclesWorld_PT_volume(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Volume"
-	bl_context = "world"
+    bl_label = "Volume"
+    bl_context = "world"
 
-	@classmethod
-	def poll(cls, context):
-		return context.world and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        return context.world and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
-		layout.active = False
+    def draw(self, context):
+        layout = self.layout
+        layout.active = False
 
-		mat = context.world
-		panel_node_draw(layout, mat, 'OUTPUT_WORLD', 'Volume')
+        mat = context.world
+        panel_node_draw(layout, mat, 'OUTPUT_WORLD', 'Volume')
 
 class CyclesMaterial_PT_surface(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Surface"
-	bl_context = "material"
+    bl_label = "Surface"
+    bl_context = "material"
 
-	@classmethod
-	def poll(cls, context):
-		return context.material and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        return context.material and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		mat = context.material
-		panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Surface')
+        mat = context.material
+        panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Surface')
 
 class CyclesMaterial_PT_volume(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Volume"
-	bl_context = "material"
+    bl_label = "Volume"
+    bl_context = "material"
 
-	@classmethod
-	def poll(cls, context):
-		return context.material and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        return context.material and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
-		layout.active = False
+    def draw(self, context):
+        layout = self.layout
+        layout.active = False
 
-		mat = context.material
-		panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Volume')
+        mat = context.material
+        panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Volume')
 
 class CyclesMaterial_PT_displacement(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Displacement"
-	bl_context = "material"
+    bl_label = "Displacement"
+    bl_context = "material"
 
-	@classmethod
-	def poll(cls, context):
-		return context.material and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        return context.material and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		mat = context.material
-		panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Displacement')
+        mat = context.material
+        panel_node_draw(layout, mat, 'OUTPUT_MATERIAL', 'Displacement')
 
 class CyclesMaterial_PT_settings(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Settings"
-	bl_context = "material"
-	bl_options = {'DEFAULT_CLOSED'}
+    bl_label = "Settings"
+    bl_context = "material"
+    bl_options = {'DEFAULT_CLOSED'}
 
-	@classmethod
-	def poll(cls, context):
-		# return context.material and CyclesButtonsPanel.poll(context)
-		return False
+    @classmethod
+    def poll(cls, context):
+        # return context.material and CyclesButtonsPanel.poll(context)
+        return False
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		mat = context.material
-	
-		row = layout.row()
-		row.label(text="Light Group:")
-		row.prop(mat, "light_group", text="")
+        mat = context.material
+    
+        row = layout.row()
+        row.label(text="Light Group:")
+        row.prop(mat, "light_group", text="")
 
 class CyclesTexture_PT_context(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = ""
-	bl_context = "texture"
-	bl_options = {'HIDE_HEADER'}
-	COMPAT_ENGINES = {'CYCLES'}
-
-	def draw(self, context):
-		layout = self.layout
-
-		tex = context.texture
-		space = context.space_data
-		pin_id = space.pin_id
-		use_pin_id = space.use_pin_id;
-		user = context.texture_user
-		node = context.texture_node
-
-		if not use_pin_id or not isinstance(pin_id, bpy.types.Texture):
-			pin_id = None
-
-		if not pin_id:
-			layout.template_texture_user()
-
-		if user:
-			layout.separator()
-
-			split = layout.split(percentage=0.65)
-			col = split.column()
-
-			if pin_id:
-				col.template_ID(space, "pin_id")
-			elif user:
-				col.template_ID(user, "texture", new="texture.new")
-			
-			if tex:
-				row = split.row()
-				row.prop(tex, "use_nodes", icon="NODETREE", text="")
-				row.label()
-
-				if not tex.use_nodes:
-					split = layout.split(percentage=0.2)
-					split.label(text="Type:")
-					split.prop(tex, "type", text="")
+    bl_label = ""
+    bl_context = "texture"
+    bl_options = {'HIDE_HEADER'}
+    COMPAT_ENGINES = {'CYCLES'}
+
+    def draw(self, context):
+        layout = self.layout
+
+        tex = context.texture
+        space = context.space_data
+        pin_id = space.pin_id
+        use_pin_id = space.use_pin_id;
+        user = context.texture_user
+        node = context.texture_node
+
+        if not use_pin_id or not isinstance(pin_id, bpy.types.Texture):
+            pin_id = None
+
+        if not pin_id:
+            layout.template_texture_user()
+
+        if user:
+            layout.separator()
+
+            split = layout.split(percentage=0.65)
+            col = split.column()
+
+            if pin_id:
+                col.template_ID(space, "pin_id")
+            elif user:
+                col.template_ID(user, "texture", new="texture.new")
+            
+            if tex:
+                row = split.row()
+                row.prop(tex, "use_nodes", icon="NODETREE", text="")
+                row.label()
+
+                if not tex.use_nodes:
+                    split = layout.split(percentage=0.2)
+                    split.label(text="Type:")
+                    split.prop(tex, "type", text="")
 
 class CyclesTexture_PT_nodes(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Nodes"
-	bl_context = "texture"
+    bl_label = "Nodes"
+    bl_context = "texture"
 
-	@classmethod
-	def poll(cls, context):
-		tex = context.texture
-		return (tex and tex.use_nodes) and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        tex = context.texture
+        return (tex and tex.use_nodes) and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		tex = context.texture
-		panel_node_draw(layout, tex, 'OUTPUT_TEXTURE', 'Color')
+        tex = context.texture
+        panel_node_draw(layout, tex, 'OUTPUT_TEXTURE', 'Color')
 
 class CyclesTexture_PT_node(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Node"
-	bl_context = "texture"
+    bl_label = "Node"
+    bl_context = "texture"
 
-	@classmethod
-	def poll(cls, context):
-		node = context.texture_node
-		return node and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        node = context.texture_node
+        return node and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
+    def draw(self, context):
+        layout = self.layout
 
-		node = context.texture_node
-		ntree = node.id_data
-		layout.template_node_view(ntree, node, None)
+        node = context.texture_node
+        ntree = node.id_data
+        layout.template_node_view(ntree, node, None)
 
 class CyclesTexture_PT_mapping(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Mapping"
-	bl_context = "texture"
+    bl_label = "Mapping"
+    bl_context = "texture"
 
-	@classmethod
-	def poll(cls, context):
-		tex = context.texture
-		node = context.texture_node
-		return (node or (tex and tex.use_nodes)) and CyclesButtonsPanel.poll(context)
+    @classmethod
+    def poll(cls, context):
+        tex = context.texture
+        node = context.texture_node
+        return (node or (tex and tex.use_nodes)) and CyclesButtonsPanel.poll(context)
 
-	def draw(self, context):
-		layout = self.layout
-		layout.label("Texture coordinate mapping goes here.");
-		layout.label("Translate, rotate, scale, projection, XYZ.")
+    def draw(self, context):
+        layout = self.layout
+        layout.label("Texture coordinate mapping goes here.");
+        layout.label("Translate, rotate, scale, projection, XYZ.")
 
 class CyclesTexture_PT_color(CyclesButtonsPanel, bpy.types.Panel):
-	bl_label = "Color"
-	bl_context = "texture"
-
-	@classmethod
-	def poll(cls, context):
-		tex = context.texture
-		node = context.texture_node
-		return (node or (tex and tex.use_nodes)) and CyclesButtonsPanel.poll(context)
-
-	def draw(self, context):
-		layout = self.layout
-		layout.label("Color modification options go here.");
-		layout.label("Ramp, brightness, contrast, saturation.")
-	
+    bl_label = "Color"
+    bl_context = "texture"
+
+    @classmethod
+    def poll(cls, context):
+        tex = context.texture
+        node = context.texture_node
+        return (node or (tex and tex.use_nodes)) and CyclesButtonsPanel.poll(context)
+
+    def draw(self, context):
+        layout = self.layout
+        layout.label("Color modification options go here.");
+        layout.label("Ramp, brightness, contrast, saturation.")
+    
 def draw_device(self, context):
-	scene = context.scene
-	layout = self.layout
+    scene = context.scene
+    layout = self.layout
 
-	if scene.render.engine == "CYCLES":
-		cycles = scene.cycles
+    if scene.render.engine == "CYCLES":
+        cscene = scene.cycles
 
-		if 'cuda' in engine.available_devices():
-			layout.prop(cycles, "device")
-		if cycles.device == 'CPU' and engine.with_osl():
-			layout.prop(cycles, "shading_system")
+        if 'cuda' in engine.available_devices():
+            layout.prop(cscene, "device")
+        if cscene.device == 'CPU' and engine.with_osl():
+            layout.prop(cscene, "shading_system")
 
 def get_panels():
-	return [
-		bpy.types.RENDER_PT_render,
-		bpy.types.RENDER_PT_output,
-		bpy.types.RENDER_PT_encoding,
-		bpy.types.RENDER_PT_dimensions,
-		bpy.types.RENDER_PT_stamp,
-		bpy.types.WORLD_PT_context_world,
-		bpy.types.DATA_PT_context_mesh,
-		bpy.types.DATA_PT_context_camera,
-		bpy.types.DATA_PT_context_lamp,
-		bpy.types.DATA_PT_texture_space,
-		bpy.types.DATA_PT_curve_texture_space,
-		bpy.types.DATA_PT_mball_texture_space,
-		bpy.types.DATA_PT_vertex_groups,
-		bpy.types.DATA_PT_shape_keys,
-		bpy.types.DATA_PT_uv_texture,
-		bpy.types.DATA_PT_vertex_colors,
-		bpy.types.DATA_PT_camera,
-		bpy.types.DATA_PT_camera_display,
-		bpy.types.DATA_PT_custom_props_mesh,
-		bpy.types.DATA_PT_custom_props_camera,
-		bpy.types.DATA_PT_custom_props_lamp,
-		bpy.types.TEXTURE_PT_clouds,
-		bpy.types.TEXTURE_PT_wood,
-		bpy.types.TEXTURE_PT_marble,
-		bpy.types.TEXTURE_PT_magic,
-		bpy.types.TEXTURE_PT_blend,
-		bpy.types.TEXTURE_PT_stucci,
-		bpy.types.TEXTURE_PT_image,
-		bpy.types.TEXTURE_PT_image_sampling,
-		bpy.types.TEXTURE_PT_image_mapping,
-		bpy.types.TEXTURE_PT_musgrave,
-		bpy.types.TEXTURE_PT_voronoi,
-		bpy.types.TEXTURE_PT_distortednoise,
-		bpy.types.TEXTURE_PT_voxeldata,
-		bpy.types.TEXTURE_PT_pointdensity,
-		bpy.types.TEXTURE_PT_pointdensity_turbulence]
+    return [
+        bpy.types.RENDER_PT_render,
+        bpy.types.RENDER_PT_output,
+        bpy.types.RENDER_PT_encoding,
+        bpy.types.RENDER_PT_dimensions,
+        bpy.types.RENDER_PT_stamp,
+        bpy.types.WORLD_PT_context_world,
+        bpy.types.DATA_PT_context_mesh,
+        bpy.types.DATA_PT_context_camera,
+        bpy.types.DATA_PT_context_lamp,
+        bpy.types.DATA_PT_texture_space,
+        bpy.types.DATA_PT_curve_texture_space,
+        bpy.types.DATA_PT_mball_texture_space,
+        bpy.types.DATA_PT_vertex_groups,
+        bpy.types.DATA_PT_shape_keys,
+        bpy.types.DATA_PT_uv_texture,
+        bpy.types.DATA_PT_vertex_colors,
+        bpy.types.DATA_PT_camera,
+        bpy.types.DATA_PT_camera_display,
+        bpy.types.DATA_PT_custom_props_mesh,
+        bpy.types.DATA_PT_custom_props_camera,
+        bpy.types.DATA_PT_custom_props_lamp,
+        bpy.types.TEXTURE_PT_clouds,
+        bpy.types.TEXTURE_PT_wood,
+        bpy.types.TEXTURE_PT_marble,
+        bpy.types.TEXTURE_PT_magic,
+        bpy.types.TEXTURE_PT_blend,
+        bpy.types.TEXTURE_PT_stucci,
+        bpy.types.TEXTURE_PT_image,
+        bpy.types.TEXTURE_PT_image_sampling,
+        bpy.types.TEXTURE_PT_image_mapping,
+        bpy.types.TEXTURE_PT_musgrave,
+        bpy.types.TEXTURE_PT_voronoi,
+        bpy.types.TEXTURE_PT_distortednoise,
+        bpy.types.TEXTURE_PT_voxeldata,
+        bpy.types.TEXTURE_PT_pointdensity,
+        bpy.types.TEXTURE_PT_pointdensity_turbulence]
 
 def register():
-	bpy.types.RENDER_PT_render.append(draw_device)
+    bpy.types.RENDER_PT_render.append(draw_device)
 
-	for panel in get_panels():
-		panel.COMPAT_ENGINES.add('CYCLES')
-	
+    for panel in get_panels():
+        panel.COMPAT_ENGINES.add('CYCLES')
+    
 def unregister():
-	bpy.types.RENDER_PT_render.remove(draw_device)
+    bpy.types.RENDER_PT_render.remove(draw_device)
 
-	for panel in get_panels():
-		panel.COMPAT_ENGINES.remove('CYCLES')
+    for panel in get_panels():
+        panel.COMPAT_ENGINES.remove('CYCLES')
 
diff --git a/intern/cycles/blender/addon/xml.py b/intern/cycles/blender/addon/xml.py
index f489f099e8b..3713da09235 100644
--- a/intern/cycles/blender/addon/xml.py
+++ b/intern/cycles/blender/addon/xml.py
@@ -25,75 +25,75 @@ import xml.etree.ElementTree as etree
 import xml.dom.minidom as dom
 
 def strip(root):
-	root.text = None
-	root.tail = None
+    root.text = None
+    root.tail = None
 
-	for elem in root:
-		strip(elem)
+    for elem in root:
+        strip(elem)
 
 def write(node, fname):
-	strip(node)
+    strip(node)
 
-	s = etree.tostring(node)
-	s = dom.parseString(s).toprettyxml()
+    s = etree.tostring(node)
+    s = dom.parseString(s).toprettyxml()
 
-	f = open(fname, "w")
-	f.write(s)
+    f = open(fname, "w")
+    f.write(s)
 
 class ExportCyclesXML(bpy.types.Operator, ExportHelper):
-	''''''
-	bl_idname = "export_mesh.cycles_xml"
-	bl_label = "Export Cycles XML"
+    ''''''
+    bl_idname = "export_mesh.cycles_xml"
+    bl_label = "Export Cycles XML"
 
-	filename_ext = ".xml"
+    filename_ext = ".xml"
 
-	@classmethod
-	def poll(cls, context):
-		return context.active_object != None
+    @classmethod
+    def poll(cls, context):
+        return context.active_object != None
 
-	def execute(self, context):
-		filepath = bpy.path.ensure_ext(self.filepath, ".xml")
+    def execute(self, context):
+        filepath = bpy.path.ensure_ext(self.filepath, ".xml")
 
-		# get mesh
-		scene = context.scene
-		object = context.object
+        # get mesh
+        scene = context.scene
+        object = context.object
 
-		if not object:
-			raise Exception("No active object")
+        if not object:
+            raise Exception("No active object")
 
-		mesh = object.to_mesh(scene, True, 'PREVIEW')
+        mesh = object.to_mesh(scene, True, 'PREVIEW')
 
-		if not mesh:
-			raise Exception("No mesh data in active object")
+        if not mesh:
+            raise Exception("No mesh data in active object")
 
-		# generate mesh node
-		nverts = ""
-		verts = ""
-		P = ""
+        # generate mesh node
+        nverts = ""
+        verts = ""
+        P = ""
 
-		for v in mesh.vertices:
-			P += "%f %f %f  " % (v.co[0], v.co[1], v.co[2])
+        for v in mesh.vertices:
+            P += "%f %f %f  " % (v.co[0], v.co[1], v.co[2])
 
-		for i, f in enumerate(mesh.faces):
-			nverts += str(len(f.vertices)) + " "
+        for i, f in enumerate(mesh.faces):
+            nverts += str(len(f.vertices)) + " "
 
-			for v in f.vertices:
-				verts += str(v) + " "
-			verts += " "
+            for v in f.vertices:
+                verts += str(v) + " "
+            verts += " "
 
-		node = etree.Element('mesh', attrib={'nverts': nverts, 'verts': verts, 'P': P})
-		
-		# write to file
-		write(node, filepath)
+        node = etree.Element('mesh', attrib={'nverts': nverts, 'verts': verts, 'P': P})
+        
+        # write to file
+        write(node, filepath)
 
-		return {'FINISHED'}
+        return {'FINISHED'}
 
 def register():
-	pass
+    pass
 
 def unregister():
-	pass
+    pass
 
 if __name__ == "__main__":
-	register()
+    register()
 
diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp
index 4223998c1f6..7be15ca0e3c 100644
--- a/intern/cycles/blender/blender_session.cpp
+++ b/intern/cycles/blender/blender_session.cpp
@@ -21,6 +21,7 @@
 #include "camera.h"
 #include "device.h"
 #include "integrator.h"
+#include "film.h"
 #include "light.h"
 #include "scene.h"
 #include "session.h"
@@ -94,7 +95,7 @@ void BlenderSession::create_session()
 	session->progress.set_cancel_callback(function_bind(&BlenderSession::test_cancel, this));
 
 	/* start rendering */
-	session->reset(width, height);
+	session->reset(width, height, session_params.passes);
 	session->start();
 }
 
@@ -118,30 +119,16 @@ void BlenderSession::render()
 void BlenderSession::write_render_result()
 {
 	/* get result */
-	DisplayBuffer *display = session->display;
-	Device *device = session->device;
-
-	if(!display->rgba.device_pointer)
-		return;
-
-	/* todo: get float buffer */
-	device->pixels_copy_from(display->rgba, 0, width, height);
-	uchar4 *rgba = (uchar4*)display->rgba.data_pointer;
-
-	vector<float4> buffer(width*height);
-	float fac = 1.0f/255.0f;
-	bool color_management = b_scene.render().use_color_management();
-
-	/* normalize */
-	for(int i = width*height - 1; i >= 0; i--) {
-		uchar4 f = rgba[i];
-		float3 rgb = make_float3(f.x, f.y, f.z)*fac;
+	RenderBuffers *buffers = session->buffers;
+	float exposure = scene->film->exposure;
+	double total_time, pass_time;
+	int pass;
+	session->progress.get_pass(pass, total_time, pass_time);
 
-		if(color_management)
-			rgb = color_srgb_to_scene_linear(rgb);
+	float4 *pixels = buffers->copy_from_device(exposure, pass);
 
-		buffer[i] = make_float4(rgb.x, rgb.y, rgb.z, 1.0f);
-	}
+	if(!pixels)
+		return;
 
 	struct RenderResult *rrp = RE_engine_begin_result((RenderEngine*)b_engine.ptr.data, 0, 0, width, height);
 	PointerRNA rrptr;
@@ -150,9 +137,11 @@ void BlenderSession::write_render_result()
 
 	BL::RenderResult::layers_iterator layer;
 	rr.layers.begin(layer);
-	rna_RenderLayer_rect_set(&layer->ptr, (float*)&buffer[0]);
+	rna_RenderLayer_rect_set(&layer->ptr, (float*)pixels);
 
 	RE_engine_end_result((RenderEngine*)b_engine.ptr.data, rrp);
+
+	delete [] pixels;
 }
 
 void BlenderSession::synchronize()
@@ -168,6 +157,9 @@ void BlenderSession::synchronize()
 		return;
 	}
 
+	/* increase passes, but never decrease */
+	session->set_passes(session_params.passes);
+
 	/* copy recalc flags, outside of mutex so we can decide to do the real
 	   synchronization at a later time to not block on running updates */
 	sync->sync_recalc();
@@ -188,7 +180,7 @@ void BlenderSession::synchronize()
 
 	/* reset if needed */
 	if(scene->need_reset())
-		session->reset(width, height);
+		session->reset(width, height, session_params.passes);
 
 	/* unlock */
 	session->scene->mutex.unlock();
@@ -225,8 +217,10 @@ bool BlenderSession::draw(int w, int h)
 		}
 
 		/* reset if requested */
-		if(reset)
-			session->reset(width, height);
+		if(reset) {
+			SessionParams session_params = BlenderSync::get_session_params(b_scene, background);
+			session->reset(width, height, session_params.passes);
+		}
 	}
 
 	/* update status and progress for 3d view draw */
diff --git a/intern/cycles/blender/blender_shader.cpp b/intern/cycles/blender/blender_shader.cpp
index db65a7f129e..9024de092b7 100644
--- a/intern/cycles/blender/blender_shader.cpp
+++ b/intern/cycles/blender/blender_shader.cpp
@@ -208,6 +208,10 @@ static ShaderNode *add_node(BL::BlendData b_data, ShaderGraph *graph, BL::Node *
 			node = new BackgroundNode();
 			break;
 		}
+		case BL::ShaderNode::type_HOLDOUT: {
+			node = new HoldoutNode();
+			break;
+		}
 		case BL::ShaderNode::type_BSDF_ANISOTROPIC: {
 			node = new WardBsdfNode();
 			break;
@@ -594,6 +598,9 @@ void BlenderSync::sync_world()
 		shader->tag_update(scene);
 	}
 
+	PointerRNA cscene = RNA_pointer_get(&b_scene.ptr, "cycles");
+	background->transparent = get_boolean(cscene, "transparent");
+
 	if(background->modified(prevbackground))
 		background->tag_update(scene);
 
diff --git a/intern/cycles/blender/blender_sync.cpp b/intern/cycles/blender/blender_sync.cpp
index a6eeb5ec1ae..ccfc3bb2bf1 100644
--- a/intern/cycles/blender/blender_sync.cpp
+++ b/intern/cycles/blender/blender_sync.cpp
@@ -18,6 +18,7 @@
 
 #include "background.h"
 #include "film.h"
+#include "../render/filter.h"
 #include "graph.h"
 #include "integrator.h"
 #include "light.h"
@@ -149,10 +150,18 @@ void BlenderSync::sync_film()
 	Film prevfilm = *film;
 
 	film->exposure = get_float(cscene, "exposure");
-	film->response = get_enum_identifier(cscene, "response_curve");
 
 	if(film->modified(prevfilm))
 		film->tag_update(scene);
+
+	Filter *filter = scene->filter;
+	Filter prevfilter = *filter;
+
+	filter->filter_type = (FilterType)RNA_enum_get(&cscene, "filter_type");
+	filter->filter_width = (filter->filter_type == FILTER_BOX)? 1.0f: get_float(cscene, "filter_width");
+
+	if(filter->modified(prevfilter))
+		filter->tag_update(scene);
 }
 
 /* Scene Parameters */
@@ -190,11 +199,22 @@ SessionParams BlenderSync::get_session_params(BL::Scene b_scene, bool background
 	foreach(DeviceType dt, types)
 		if(dt == dtype)
 			params.device_type = dtype;
+			
+	/* Background */
+	params.background = background;
+			
+	/* passes */
+	if(background) {
+		params.passes = get_int(cscene, "passes");
+	}
+	else {
+		params.passes = get_int(cscene, "preview_passes");
+		if(params.passes == 0)
+			params.passes = INT_MAX;
+	}
 
 	/* other parameters */
 	params.threads = b_scene.render().threads();
-	params.background = background;
-	params.passes = (background)? get_int(cscene, "passes"): INT_MAX;
 	params.tile_size = get_int(cscene, "debug_tile_size");
 	params.min_size = get_int(cscene, "debug_min_size");
 	params.cancel_timeout = get_float(cscene, "debug_cancel_timeout");
diff --git a/intern/cycles/cmake/external_libs.cmake b/intern/cycles/cmake/external_libs.cmake
index b23e3b292c5..d53be9f6fe6 100644
--- a/intern/cycles/cmake/external_libs.cmake
+++ b/intern/cycles/cmake/external_libs.cmake
@@ -94,18 +94,14 @@ if(WITH_CYCLES_CUDA)
 	set(CYCLES_CUDA_ARCH sm_10 sm_11 sm_12 sm_13 sm_20 sm_21 CACHE STRING "CUDA architectures to build for")
 	set(CYCLES_CUDA_MAXREG 24 CACHE STRING "CUDA maximum number of register to use")
 
-	find_path(CUDA_INCLUDES cuda.h ${CYCLES_CUDA}/include NO_DEFAULT_PATH)
 	find_program(CUDA_NVCC NAMES nvcc PATHS ${CYCLES_CUDA}/bin NO_DEFAULT_PATH)
 
-	if(CUDA_INCLUDES AND CUDA_NVCC)
-		message(STATUS "CUDA includes = ${CUDA_INCLUDES}")
+	if(CUDA_NVCC)
 		message(STATUS "CUDA nvcc = ${CUDA_NVCC}")
 	else()
-		message(STATUS "CUDA not found")
+		message(STATUS "CUDA compiler not found")
 	endif()
 
-	include_directories(${CUDA_INCLUDES})
-
 endif()
 
 ###########################################################################
diff --git a/intern/cycles/device/device.cpp b/intern/cycles/device/device.cpp
index cd0d64f37dd..2f2d34dcef2 100644
--- a/intern/cycles/device/device.cpp
+++ b/intern/cycles/device/device.cpp
@@ -86,10 +86,15 @@ void Device::pixels_free(device_memory& mem)
 	mem_free(mem);
 }
 
-void Device::draw_pixels(device_memory& rgba, int y, int w, int h, int width, int height)
+void Device::draw_pixels(device_memory& rgba, int y, int w, int h, int width, int height, bool transparent)
 {
 	pixels_copy_from(rgba, y, w, h);
 
+	if(transparent) {
+		glEnable(GL_BLEND);
+		glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
+	}
+
 	glPixelZoom((float)width/(float)w, (float)height/(float)h);
 	glRasterPos2f(0, y);
 
@@ -97,6 +102,9 @@ void Device::draw_pixels(device_memory& rgba, int y, int w, int h, int width, in
 
 	glRasterPos2f(0.0f, 0.0f);
 	glPixelZoom(1.0f, 1.0f);
+
+	if(transparent)
+		glDisable(GL_BLEND);
 }
 
 Device *Device::create(DeviceType type, bool background, int threads)
diff --git a/intern/cycles/device/device.h b/intern/cycles/device/device.h
index 8006eea1a5f..87683ae4719 100644
--- a/intern/cycles/device/device.h
+++ b/intern/cycles/device/device.h
@@ -115,7 +115,7 @@ public:
 	
 	/* opengl drawing */
 	virtual void draw_pixels(device_memory& mem, int y, int w, int h,
-		int width, int height);
+		int width, int height, bool transparent);
 
 #ifdef WITH_NETWORK
 	/* networking */
diff --git a/intern/cycles/device/device_cuda.cpp b/intern/cycles/device/device_cuda.cpp
index 0537e231f44..cc1c1275bf4 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -607,7 +607,7 @@ public:
 		}
 	}
 
-	void draw_pixels(device_memory& mem, int y, int w, int h, int width, int height)
+	void draw_pixels(device_memory& mem, int y, int w, int h, int width, int height, bool transparent)
 	{
 		if(!background) {
 			PixelMem pmem = pixel_mem_map[mem.device_pointer];
@@ -621,11 +621,16 @@ public:
 			
 			glEnable(GL_TEXTURE_2D);
 			
+			if(transparent) {
+				glEnable(GL_BLEND);
+				glBlendFunc(GL_ONE, GL_ONE_MINUS_SRC_ALPHA);
+			}
+
 			glColor3f(1.0f, 1.0f, 1.0f);
 
 			glPushMatrix();
 			glTranslatef(0.0f, (float)y, 0.0f);
-			
+				
 			glBegin(GL_QUADS);
 			
 			glTexCoord2f(0.0f, 0.0f);
@@ -640,6 +645,9 @@ public:
 			glEnd();
 
 			glPopMatrix();
+
+			if(transparent)
+				glDisable(GL_BLEND);
 			
 			glBindTexture(GL_TEXTURE_2D, 0);
 			glDisable(GL_TEXTURE_2D);
@@ -649,7 +657,7 @@ public:
 			return;
 		}
 
-		Device::draw_pixels(mem, y, w, h, width, height);
+		Device::draw_pixels(mem, y, w, h, width, height, transparent);
 	}
 
 	void task_add(DeviceTask& task)
diff --git a/intern/cycles/device/device_multi.cpp b/intern/cycles/device/device_multi.cpp
index e48df93737d..cae63596d3b 100644
--- a/intern/cycles/device/device_multi.cpp
+++ b/intern/cycles/device/device_multi.cpp
@@ -244,7 +244,7 @@ public:
 		mem.device_pointer = tmp;
 	}
 
-	void draw_pixels(device_memory& rgba, int y, int w, int h, int width, int height)
+	void draw_pixels(device_memory& rgba, int y, int w, int h, int width, int height, bool transparent)
 	{
 		device_ptr tmp = rgba.device_pointer;
 		int i = 0, sub_h = h/devices.size();
@@ -255,7 +255,7 @@ public:
 			/* adjust math for w/width */
 
 			rgba.device_pointer = sub.ptr_map[tmp];
-			sub.device->draw_pixels(rgba, sy, w, sh, width, height);
+			sub.device->draw_pixels(rgba, sy, w, sh, width, height, transparent);
 			i++;
 		}
 
diff --git a/intern/cycles/kernel/kernel.cpp b/intern/cycles/kernel/kernel.cpp
index 656e90eb526..d9ba237d964 100644
--- a/intern/cycles/kernel/kernel.cpp
+++ b/intern/cycles/kernel/kernel.cpp
@@ -129,18 +129,6 @@ void kernel_tex_copy(KernelGlobals *kg, const char *name, device_ptr mem, size_t
 		kg->__filter_table.data = (float*)mem;
 		kg->__filter_table.width = width;
 	}
-	else if(strcmp(name, "__response_curve_R") == 0) {
-		kg->__response_curve_R.data = (float*)mem;
-		kg->__response_curve_R.width = width;
-	}
-	else if(strcmp(name, "__response_curve_B") == 0) {
-		kg->__response_curve_B.data = (float*)mem;
-		kg->__response_curve_B.width = width;
-	}
-	else if(strcmp(name, "__response_curve_G") == 0) {
-		kg->__response_curve_G.data = (float*)mem;
-		kg->__response_curve_G.width = width;
-	}
 	else if(strcmp(name, "__sobol_directions") == 0) {
 		kg->__sobol_directions.data = (uint*)mem;
 		kg->__sobol_directions.width = width;
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index b7b29d46323..3a0eff5210c 100644
--- a/intern/cycles/kernel/kernel_compat_cuda.h
+++ b/intern/cycles/kernel/kernel_compat_cuda.h
@@ -33,6 +33,7 @@ CCL_NAMESPACE_BEGIN
 
 #define __device  __device__ __inline__
 #define __device_inline  __device__ __inline__
+#define __device_noinline  __device__ __noinline__
 #define __global
 #define __shared __shared__
 #define __constant
diff --git a/intern/cycles/kernel/kernel_compat_opencl.h b/intern/cycles/kernel/kernel_compat_opencl.h
index 3d493c61fe4..287bf320881 100644
--- a/intern/cycles/kernel/kernel_compat_opencl.h
+++ b/intern/cycles/kernel/kernel_compat_opencl.h
@@ -30,6 +30,7 @@
 /* in opencl all functions are device functions, so leave this empty */
 #define __device
 #define __device_inline
+#define __device_noinline
 
 /* no assert in opencl */
 #define kernel_assert(cond)
diff --git a/intern/cycles/kernel/kernel_film.h b/intern/cycles/kernel/kernel_film.h
index 0bbd6c202a1..f6351a73295 100644
--- a/intern/cycles/kernel/kernel_film.h
+++ b/intern/cycles/kernel/kernel_film.h
@@ -20,21 +20,17 @@ CCL_NAMESPACE_BEGIN
 
 __device float4 film_map(KernelGlobals *kg, float4 irradiance, int pass)
 {
-	float scale = kernel_data.film.exposure*(1.0f/(pass+1));
+	float scale = 1.0f/(float)(pass+1);
+	float exposure = kernel_data.film.exposure;
 	float4 result = irradiance*scale;
 
-	if(kernel_data.film.use_response_curve) {
-		/* camera response curve */
-		result.x = kernel_tex_interp(__response_curve_R, result.x);
-		result.y = kernel_tex_interp(__response_curve_G, result.y);
-		result.z = kernel_tex_interp(__response_curve_B, result.z);
-	}
-	else {
-		/* conversion to srgb */
-		result.x = color_scene_linear_to_srgb(result.x);
-		result.y = color_scene_linear_to_srgb(result.y);
-		result.z = color_scene_linear_to_srgb(result.z);
-	}
+	/* conversion to srgb */
+	result.x = color_scene_linear_to_srgb(result.x*exposure);
+	result.y = color_scene_linear_to_srgb(result.y*exposure);
+	result.z = color_scene_linear_to_srgb(result.z*exposure);
+
+	/* clamp since alpha might be > 1.0 due to russian roulette */
+	result.w = clamp(result.w, 0.0f, 1.0f);
 
 	return result;
 }
@@ -47,7 +43,7 @@ __device uchar4 film_float_to_byte(float4 color)
 	result.x = (uchar)clamp(color.x*255.0f, 0.0f, 255.0f);
 	result.y = (uchar)clamp(color.y*255.0f, 0.0f, 255.0f);
 	result.z = (uchar)clamp(color.z*255.0f, 0.0f, 255.0f);
-	result.w = 255;
+	result.w = (uchar)clamp(color.w*255.0f, 0.0f, 255.0f);
 
 	return result;
 }
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index 47aeea4f5c5..2cbb4ae306e 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -105,10 +105,11 @@ __device int path_flag_from_label(int path_flag, int label)
 	return path_flag;
 }
 
-__device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray ray, float3 throughput)
+__device float4 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray ray, float3 throughput)
 {
 	/* initialize */
 	float3 L = make_float3(0.0f, 0.0f, 0.0f);
+	float Ltransparent = 0.0f;
 
 #ifdef __EMISSION__
 	float ray_pdf = 0.0f;
@@ -123,14 +124,20 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 
 		if(!scene_intersect(kg, &ray, false, &isect)) {
 			/* eval background shader if nothing hit */
+			if(kernel_data.background.transparent && (path_flag & PATH_RAY_CAMERA)) {
+				Ltransparent += average(throughput);
+			}
+			else {
 #ifdef __BACKGROUND__
-			ShaderData sd;
-			shader_setup_from_background(kg, &sd, &ray);
-			L += throughput*shader_eval_background(kg, &sd, path_flag);
-			shader_release(kg, &sd);
+				ShaderData sd;
+				shader_setup_from_background(kg, &sd, &ray);
+				L += throughput*shader_eval_background(kg, &sd, path_flag);
+				shader_release(kg, &sd);
 #else
-			L += make_float3(0.8f, 0.8f, 0.8f);
+				L += make_float3(0.8f, 0.8f, 0.8f);
 #endif
+			}
+
 			break;
 		}
 
@@ -140,6 +147,22 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 		float rbsdf = path_rng(kg, rng, pass, rng_offset + PRNG_BSDF);
 		shader_eval_surface(kg, &sd, rbsdf, path_flag);
 
+#ifdef __HOLDOUT__
+		if((sd.flag & SD_HOLDOUT) && (path_flag & PATH_RAY_CAMERA)) {
+			float3 holdout_weight = shader_holdout_eval(kg, &sd);
+
+			if(kernel_data.background.transparent) {
+				Ltransparent += average(holdout_weight*throughput);
+			}
+			else {
+				ShaderData sd;
+				shader_setup_from_background(kg, &sd, &ray);
+				L += holdout_weight*throughput*shader_eval_background(kg, &sd, path_flag);
+				shader_release(kg, &sd);
+			}
+		}
+#endif
+
 #ifdef __EMISSION__
 		/* emission */
 		if(kernel_data.integrator.use_emission) {
@@ -166,6 +189,12 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 		}
 #endif
 
+		/* no BSDF? we can stop here */
+		if(!(sd.flag & SD_BSDF)) {
+			path_flag &= ~PATH_RAY_CAMERA;
+			break;
+		}
+
 		/* sample BSDF */
 		float bsdf_pdf;
 		float3 bsdf_eval;
@@ -180,8 +209,10 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 
 		shader_release(kg, &sd);
 
-		if(bsdf_pdf == 0.0f || is_zero(bsdf_eval))
+		if(bsdf_pdf == 0.0f || is_zero(bsdf_eval)) {
+			path_flag &= ~PATH_RAY_CAMERA;
 			break;
+		}
 
 		/* modify throughput */
 		throughput *= bsdf_eval/bsdf_pdf;
@@ -197,8 +228,10 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 		float probability = path_terminate_probability(kg, bounce, throughput);
 		float terminate = path_rng(kg, rng, pass, rng_offset + PRNG_TERMINATE);
 
-		if(terminate >= probability)
+		if(terminate >= probability) {
+			path_flag &= ~PATH_RAY_CAMERA;
 			break;
+		}
 
 		throughput /= probability;
 
@@ -212,7 +245,7 @@ __device float3 kernel_path_integrate(KernelGlobals *kg, RNG *rng, int pass, Ray
 #endif
 	}
 
-	return L;
+	return make_float4(L.x, L.y, L.z, 1.0f - Ltransparent);
 }
 
 __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __global uint *rng_state, int pass, int x, int y)
@@ -236,25 +269,19 @@ __device void kernel_path_trace(KernelGlobals *kg, __global float4 *buffer, __gl
 	/* integrate */
 #ifdef __MODIFY_TP__
 	float3 throughput = path_terminate_modified_throughput(kg, buffer, x, y, pass);
-	float3 L = kernel_path_integrate(kg, &rng, pass, ray, throughput)/throughput;
+	float4 L = kernel_path_integrate(kg, &rng, pass, ray, throughput)/throughput;
 #else
 	float3 throughput = make_float3(1.0f, 1.0f, 1.0f);
-	float3 L = kernel_path_integrate(kg, &rng, pass, ray, throughput);
+	float4 L = kernel_path_integrate(kg, &rng, pass, ray, throughput);
 #endif
 
 	/* accumulate result in output buffer */
 	int index = x + y*kernel_data.cam.width;
 
-	float4 result;
-	result.x = L.x;
-	result.y = L.y;
-	result.z = L.z;
-	result.w = 1.0f;
-
 	if(pass == 0)
-		buffer[index] = result;
+		buffer[index] = L;
 	else
-		buffer[index] += result;
+		buffer[index] += L;
 
 	path_rng_end(kg, rng_state, rng, x, y);
 }
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index c4fc65596b2..ed3aec21207 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -342,6 +342,26 @@ __device void shader_emissive_sample(KernelGlobals *kg, ShaderData *sd,
 	}
 }
 
+/* Holdout */
+
+__device float3 shader_holdout_eval(KernelGlobals *kg, ShaderData *sd)
+{
+#ifdef WITH_OSL
+	if(kg->osl.use) {
+		return OSLShader::holdout_eval(sd);
+	}
+	else
+#endif
+	{
+#ifdef __SVM__
+		if(sd->svm_closure == CLOSURE_HOLDOUT_ID)
+			return make_float3(1.0f, 1.0f, 1.0f);
+		else
+			return make_float3(0.0f, 0.0f, 0.0f);
+#endif
+	}
+}
+
 /* Surface Evaluation */
 
 __device void shader_eval_surface(KernelGlobals *kg, ShaderData *sd,
diff --git a/intern/cycles/kernel/kernel_textures.h b/intern/cycles/kernel/kernel_textures.h
index bd44ed7eee5..dd4566ec184 100644
--- a/intern/cycles/kernel/kernel_textures.h
+++ b/intern/cycles/kernel/kernel_textures.h
@@ -38,9 +38,6 @@ KERNEL_TEX(uint4, texture_uint4, __svm_nodes)
 
 /* camera/film */
 KERNEL_TEX(float, texture_float, __filter_table)
-KERNEL_TEX(float, texture_float, __response_curve_R)
-KERNEL_TEX(float, texture_float, __response_curve_G)
-KERNEL_TEX(float, texture_float, __response_curve_B)
 
 /* sobol */
 KERNEL_TEX(uint, texture_uint, __sobol_directions)
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index c0cb3fc8a09..06581d593e4 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -37,6 +37,7 @@ CCL_NAMESPACE_BEGIN
 #ifndef __KERNEL_OPENCL__
 #define __SVM__
 #define __TEXTURES__
+#define __HOLDOUT__
 #endif
 #define __RAY_DIFFERENTIALS__
 #define __CAMERA_CLIPPING__
@@ -194,8 +195,10 @@ struct FlatClosure {
 enum ShaderDataFlag {
 	SD_BACKFACING = 1,		/* backside of surface? */
 	SD_EMISSION = 2,		/* have emissive closure? */
-	SD_BSDF_HAS_EVAL = 4,	/* have non-singular bsdf closure? */
-	SD_BSDF_GLOSSY = 8		/* have glossy bsdf */
+	SD_BSDF = 4,			/* have bsdf closure? */
+	SD_BSDF_HAS_EVAL = 8,	/* have non-singular bsdf closure? */
+	SD_BSDF_GLOSSY = 16,	/* have glossy bsdf */
+	SD_HOLDOUT = 32			/* have holdout closure? */
 };
 
 typedef struct ShaderData {
@@ -257,6 +260,8 @@ typedef struct ShaderData {
 		float emissive_sample_sum;
 		float volume_sample_sum;
 
+		float3 holdout_weight;
+
 		float randb;
 	} osl_closure;
 
@@ -313,14 +318,14 @@ typedef struct KernelCamera {
 
 typedef struct KernelFilm {
 	float exposure;
-	int use_response_curve;
-	int pad1, pad2;
+	int pad1, pad2, pad3;
 } KernelFilm;
 
 typedef struct KernelBackground {
 	/* only shader index */
 	int shader;
-	int pad1, pad2, pad3;
+	int transparent;
+	int pad1, pad2;
 } KernelBackground;
 
 typedef struct KernelSunSky {
diff --git a/intern/cycles/kernel/osl/osl_shader.cpp b/intern/cycles/kernel/osl/osl_shader.cpp
index a86946a680e..007d14b526b 100644
--- a/intern/cycles/kernel/osl/osl_shader.cpp
+++ b/intern/cycles/kernel/osl/osl_shader.cpp
@@ -145,9 +145,11 @@ static void flatten_surface_closure_tree(ShaderData *sd, bool no_glossy,
 
 					/* scattering flags */
 					if(scattering == OSL::Labels::DIFFUSE)
-						sd->flag |= SD_BSDF_HAS_EVAL;
+						sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 					else if(scattering == OSL::Labels::GLOSSY)
-						sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+						sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+					else
+						sd->flag |= SD_BSDF;
 
 					/* add */
 					sd->osl_closure.bsdf[sd->osl_closure.num_bsdf++] = flat;
@@ -170,8 +172,11 @@ static void flatten_surface_closure_tree(ShaderData *sd, bool no_glossy,
 					sd->osl_closure.emissive[sd->osl_closure.num_emissive++] = flat;
 					break;
 				}
-				case ClosurePrimitive::BSSRDF:
 				case ClosurePrimitive::Holdout:
+					sd->osl_closure.holdout_weight += weight;
+					sd->flag |= SD_HOLDOUT;
+					break;
+				case ClosurePrimitive::BSSRDF:
 				case ClosurePrimitive::Debug:
 					break; /* not implemented */
 				case ClosurePrimitive::Background:
@@ -212,6 +217,7 @@ void OSLShader::eval_surface(KernelGlobals *kg, ShaderData *sd, float randb, int
 	sd->osl_closure.emissive_sample_sum = 0.0f;
 	sd->osl_closure.num_bsdf = 0;
 	sd->osl_closure.num_emissive = 0;
+	sd->osl_closure.holdout_weight = make_float3(0.0f, 0.0f, 0.0f);
 	sd->osl_closure.randb = randb;
 
 	if(globals->Ci) {
@@ -555,5 +561,12 @@ float3 OSLShader::volume_eval_phase(const ShaderData *sd, const float3 omega_in,
 	return eval;
 }
 
+/* Holdout Closure */
+
+float3 OSLShader::holdout_eval(const ShaderData *sd)
+{
+	return sd->osl_closure.holdout_weight;
+}
+
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/kernel/osl/osl_shader.h b/intern/cycles/kernel/osl/osl_shader.h
index 2e5279c12ce..9b578b159ae 100644
--- a/intern/cycles/kernel/osl/osl_shader.h
+++ b/intern/cycles/kernel/osl/osl_shader.h
@@ -74,6 +74,7 @@ public:
 		float3 *eval, float3 *I, float *pdf);
 	static float3 volume_eval_phase(const ShaderData *sd, const float3 omega_in,
 		const float3 omega_out);
+	static float3 holdout_eval(const ShaderData *sd);
 
 	/* release */
 	static void release(KernelGlobals *kg, const ShaderData *sd);
diff --git a/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h b/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h
index 04e4ccb8313..a04f4e2b076 100644
--- a/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h
+++ b/intern/cycles/kernel/svm/bsdf_ashikhmin_velvet.h
@@ -50,7 +50,7 @@ __device void bsdf_ashikhmin_velvet_setup(ShaderData *sd, float3 N, float sigma)
 	self->m_invsigma2 = 1.0f/(sigma * sigma);
 
 	sd->svm_closure = CLOSURE_BSDF_ASHIKHMIN_VELVET_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 }
 
 __device void bsdf_ashikhmin_velvet_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_diffuse.h b/intern/cycles/kernel/svm/bsdf_diffuse.h
index 00493e72203..dcd29534109 100644
--- a/intern/cycles/kernel/svm/bsdf_diffuse.h
+++ b/intern/cycles/kernel/svm/bsdf_diffuse.h
@@ -47,7 +47,7 @@ __device void bsdf_diffuse_setup(ShaderData *sd, float3 N)
 	//self->m_N = N;
 
 	sd->svm_closure = CLOSURE_BSDF_DIFFUSE_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 }
 
 __device void bsdf_diffuse_blur(ShaderData *sd, float roughness)
@@ -110,7 +110,7 @@ __device void bsdf_translucent_setup(ShaderData *sd, float3 N)
 	//self->m_N = N;
 
 	sd->svm_closure = CLOSURE_BSDF_TRANSLUCENT_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL;
 }
 
 __device void bsdf_translucent_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_microfacet.h b/intern/cycles/kernel/svm/bsdf_microfacet.h
index a948ba06871..c98092b2f8f 100644
--- a/intern/cycles/kernel/svm/bsdf_microfacet.h
+++ b/intern/cycles/kernel/svm/bsdf_microfacet.h
@@ -58,7 +58,7 @@ __device void bsdf_microfacet_ggx_setup(ShaderData *sd, float3 N, float ag, floa
 	else
 		sd->svm_closure = CLOSURE_BSDF_MICROFACET_GGX_ID;
 
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_microfacet_ggx_blur(ShaderData *sd, float roughness)
@@ -278,7 +278,7 @@ __device void bsdf_microfacet_beckmann_setup(ShaderData *sd, float3 N, float ab,
 	else
 		sd->svm_closure = CLOSURE_BSDF_MICROFACET_BECKMANN_ID;
 
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_microfacet_beckmann_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_reflection.h b/intern/cycles/kernel/svm/bsdf_reflection.h
index 287cc9c2506..2c49eb7ef24 100644
--- a/intern/cycles/kernel/svm/bsdf_reflection.h
+++ b/intern/cycles/kernel/svm/bsdf_reflection.h
@@ -47,6 +47,7 @@ __device void bsdf_reflection_setup(ShaderData *sd, float3 N)
 	//self->m_N = N;
 
 	sd->svm_closure = CLOSURE_BSDF_REFLECTION_ID;
+	sd->flag |= SD_BSDF;
 }
 
 __device void bsdf_reflection_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_refraction.h b/intern/cycles/kernel/svm/bsdf_refraction.h
index 55a914f8334..534945f4f0b 100644
--- a/intern/cycles/kernel/svm/bsdf_refraction.h
+++ b/intern/cycles/kernel/svm/bsdf_refraction.h
@@ -48,6 +48,7 @@ __device void bsdf_refraction_setup(ShaderData *sd, float3 N, float eta)
 	self->m_eta = eta;
 
 	sd->svm_closure = CLOSURE_BSDF_REFRACTION_ID;
+	sd->flag |= SD_BSDF;
 }
 
 __device void bsdf_refraction_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_transparent.h b/intern/cycles/kernel/svm/bsdf_transparent.h
index e689e3db357..30288bf251d 100644
--- a/intern/cycles/kernel/svm/bsdf_transparent.h
+++ b/intern/cycles/kernel/svm/bsdf_transparent.h
@@ -38,6 +38,7 @@ CCL_NAMESPACE_BEGIN
 __device void bsdf_transparent_setup(ShaderData *sd)
 {
 	sd->svm_closure = CLOSURE_BSDF_TRANSPARENT_ID;
+	sd->flag |= SD_BSDF;
 }
 
 __device void bsdf_transparent_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_ward.h b/intern/cycles/kernel/svm/bsdf_ward.h
index c54418afa77..6680644a1a7 100644
--- a/intern/cycles/kernel/svm/bsdf_ward.h
+++ b/intern/cycles/kernel/svm/bsdf_ward.h
@@ -54,7 +54,7 @@ __device void bsdf_ward_setup(ShaderData *sd, float3 N, float3 T, float ax, floa
 	self->m_ay = clamp(ay, 1e-5f, 1.0f);
 
 	sd->svm_closure = CLOSURE_BSDF_WARD_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_ward_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/bsdf_westin.h b/intern/cycles/kernel/svm/bsdf_westin.h
index 6031012d0ca..55fc8286529 100644
--- a/intern/cycles/kernel/svm/bsdf_westin.h
+++ b/intern/cycles/kernel/svm/bsdf_westin.h
@@ -51,7 +51,7 @@ __device void bsdf_westin_backscatter_setup(ShaderData *sd, float3 N, float roug
 	self->m_invroughness = 1.0f/roughness;
 
 	sd->svm_closure = CLOSURE_BSDF_WESTIN_BACKSCATTER_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_westin_backscatter_blur(ShaderData *sd, float roughness)
@@ -146,7 +146,7 @@ __device void bsdf_westin_sheen_setup(ShaderData *sd, float3 N, float edginess)
 	self->m_edginess = edginess;
 
 	sd->svm_closure = CLOSURE_BSDF_WESTIN_SHEEN_ID;
-	sd->flag |= SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
+	sd->flag |= SD_BSDF|SD_BSDF_HAS_EVAL|SD_BSDF_GLOSSY;
 }
 
 __device void bsdf_westin_sheen_blur(ShaderData *sd, float roughness)
diff --git a/intern/cycles/kernel/svm/svm.h b/intern/cycles/kernel/svm/svm.h
index 68e54d9752e..3a2b6bfb1ae 100644
--- a/intern/cycles/kernel/svm/svm.h
+++ b/intern/cycles/kernel/svm/svm.h
@@ -148,7 +148,7 @@ CCL_NAMESPACE_BEGIN
 
 /* Main Interpreter Loop */
 
-__device void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, float randb, int path_flag)
+__device_noinline void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type, float randb, int path_flag)
 {
 	float stack[SVM_STACK_SIZE];
 	float closure_weight = 1.0f;
@@ -172,6 +172,8 @@ __device void svm_eval_nodes(KernelGlobals *kg, ShaderData *sd, ShaderType type,
 			svm_node_closure_emission(sd);
 		else if(node.x == NODE_CLOSURE_BACKGROUND)
 			svm_node_closure_background(sd);
+		else if(node.x == NODE_CLOSURE_HOLDOUT)
+			svm_node_closure_holdout(sd);
 		else if(node.x == NODE_CLOSURE_SET_WEIGHT)
 			svm_node_closure_set_weight(sd, node.y, node.z, node.w);
 		else if(node.x == NODE_CLOSURE_WEIGHT)
diff --git a/intern/cycles/kernel/svm/svm_closure.h b/intern/cycles/kernel/svm/svm_closure.h
index a7462e2c358..1efadcefbd5 100644
--- a/intern/cycles/kernel/svm/svm_closure.h
+++ b/intern/cycles/kernel/svm/svm_closure.h
@@ -136,6 +136,12 @@ __device void svm_node_closure_background(ShaderData *sd)
 	sd->svm_closure = CLOSURE_BACKGROUND_ID;
 }
 
+__device void svm_node_closure_holdout(ShaderData *sd)
+{
+	sd->svm_closure = CLOSURE_HOLDOUT_ID;
+	sd->flag |= SD_HOLDOUT;
+}
+
 /* Closure Nodes */
 
 __device void svm_node_closure_set_weight(ShaderData *sd, uint r, uint g, uint b)
diff --git a/intern/cycles/kernel/svm/svm_convert.h b/intern/cycles/kernel/svm/svm_convert.h
index 14925569bfb..188b0489d9e 100644
--- a/intern/cycles/kernel/svm/svm_convert.h
+++ b/intern/cycles/kernel/svm/svm_convert.h
@@ -30,7 +30,7 @@ __device void svm_node_convert(ShaderData *sd, float *stack, uint type, uint fro
 		}
 		case NODE_CONVERT_CF: {
 			float3 f = stack_load_float3(stack, from);
-			float g = f.x*0.2126f + f.y*0.7152f + f.z*0.0722f;
+			float g = linear_rgb_to_gray(f);
 			stack_store_float(stack, to, g);
 			break;
 		}
diff --git a/intern/cycles/kernel/svm/svm_types.h b/intern/cycles/kernel/svm/svm_types.h
index 5cbd05e4400..786478c0c03 100644
--- a/intern/cycles/kernel/svm/svm_types.h
+++ b/intern/cycles/kernel/svm/svm_types.h
@@ -81,7 +81,8 @@ typedef enum NodeType {
 	NODE_EMISSION_SET_WEIGHT_TOTAL = 4300,
 	NODE_ATTR_BUMP_DX = 4400,
 	NODE_ATTR_BUMP_DY = 4500,
-	NODE_TEX_ENVIRONMENT = 4600
+	NODE_TEX_ENVIRONMENT = 4600,
+	NODE_CLOSURE_HOLDOUT = 4700
 } NodeType;
 
 typedef enum NodeAttributeType {
diff --git a/intern/cycles/render/CMakeLists.txt b/intern/cycles/render/CMakeLists.txt
index c3d5ae1ea05..14859aa4414 100644
--- a/intern/cycles/render/CMakeLists.txt
+++ b/intern/cycles/render/CMakeLists.txt
@@ -15,7 +15,6 @@ set(sources
 	buffers.cpp
 	camera.cpp
 	film.cpp
-	film_response.cpp
 	filter.cpp
 	graph.cpp
 	image.cpp
@@ -39,7 +38,6 @@ set(headers
 	buffers.h
 	camera.h
 	film.h
-	film_response.h
 	filter.h
 	graph.h
 	image.h
diff --git a/intern/cycles/render/background.cpp b/intern/cycles/render/background.cpp
index 871fe128514..bf4131b6d38 100644
--- a/intern/cycles/render/background.cpp
+++ b/intern/cycles/render/background.cpp
@@ -31,6 +31,7 @@ CCL_NAMESPACE_BEGIN
 
 Background::Background()
 {
+	transparent = false;
 	need_update = true;
 }
 
@@ -45,8 +46,9 @@ void Background::device_update(Device *device, DeviceScene *dscene, Scene *scene
 	
 	device_free(device, dscene);
 
-	/* set shader index */
+	/* set shader index and transparent option */
 	KernelBackground *kbackground = &dscene->data.background;
+	kbackground->transparent = transparent;
 	kbackground->shader = scene->shader_manager->get_shader_id(scene->default_background);
 
 	need_update = false;
@@ -58,7 +60,7 @@ void Background::device_free(Device *device, DeviceScene *dscene)
 
 bool Background::modified(const Background& background)
 {
-	return false;
+	return !(transparent == background.transparent);
 }
 
 void Background::tag_update(Scene *scene)
diff --git a/intern/cycles/render/background.h b/intern/cycles/render/background.h
index a048cbc535c..16d2ccb6caa 100644
--- a/intern/cycles/render/background.h
+++ b/intern/cycles/render/background.h
@@ -29,6 +29,7 @@ class Scene;
 
 class Background {
 public:
+	bool transparent;
 	bool need_update;
 
 	Background();
diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp
index c1d9158c85e..62dc17960ae 100644
--- a/intern/cycles/render/buffers.cpp
+++ b/intern/cycles/render/buffers.cpp
@@ -24,6 +24,7 @@
 #include "util_debug.h"
 #include "util_hash.h"
 #include "util_image.h"
+#include "util_math.h"
 #include "util_opengl.h"
 #include "util_time.h"
 #include "util_types.h"
@@ -84,6 +85,33 @@ void RenderBuffers::reset(Device *device, int width_, int height_)
 	device->mem_copy_to(rng_state);
 }
 
+float4 *RenderBuffers::copy_from_device(float exposure, int pass)
+{
+	if(!buffer.device_pointer)
+		return NULL;
+
+	device->mem_copy_from(buffer, 0, buffer.memory_size());
+
+	float4 *out = new float4[width*height];
+	float4 *in = (float4*)buffer.data_pointer;
+	float scale = 1.0f/(float)pass;
+	
+	for(int i = width*height - 1; i >= 0; i--) {
+		float4 rgba = in[i]*scale;
+
+		rgba.x = rgba.x*exposure;
+		rgba.y = rgba.y*exposure;
+		rgba.z = rgba.z*exposure;
+
+		/* clamp since alpha might be > 1.0 due to russian roulette */
+		rgba.w = clamp(rgba.w, 0.0f, 1.0f);
+
+		out[i] = rgba;
+	}
+
+	return out;
+}
+
 /* Display Buffer */
 
 DisplayBuffer::DisplayBuffer(Device *device_)
@@ -93,6 +121,7 @@ DisplayBuffer::DisplayBuffer(Device *device_)
 	height = 0;
 	draw_width = 0;
 	draw_height = 0;
+	transparent = true; /* todo: determine from background */
 }
 
 DisplayBuffer::~DisplayBuffer()
@@ -132,10 +161,36 @@ void DisplayBuffer::draw_set(int width_, int height_)
 	draw_height = height_;
 }
 
+void DisplayBuffer::draw_transparency_grid()
+{
+	GLubyte checker_stipple_sml[32*32/8] = {
+		255,0,255,0,255,0,255,0,255,0,255,0,255,0,255,0, \
+		255,0,255,0,255,0,255,0,255,0,255,0,255,0,255,0, \
+		0,255,0,255,0,255,0,255,0,255,0,255,0,255,0,255, \
+		0,255,0,255,0,255,0,255,0,255,0,255,0,255,0,255, \
+		255,0,255,0,255,0,255,0,255,0,255,0,255,0,255,0, \
+		255,0,255,0,255,0,255,0,255,0,255,0,255,0,255,0, \
+		0,255,0,255,0,255,0,255,0,255,0,255,0,255,0,255, \
+		0,255,0,255,0,255,0,255,0,255,0,255,0,255,0,255, \
+	};
+
+	glColor4ub(50, 50, 50, 255);
+	glRectf(0, 0, width, height);
+	glEnable(GL_POLYGON_STIPPLE);
+	glColor4ub(55, 55, 55, 255);
+	glPolygonStipple(checker_stipple_sml);
+	glRectf(0, 0, width, height);
+	glDisable(GL_POLYGON_STIPPLE);
+}
+
 void DisplayBuffer::draw(Device *device)
 {
-	if(draw_width != 0 && draw_height != 0)
-		device->draw_pixels(rgba, 0, draw_width, draw_height, width, height);
+	if(draw_width != 0 && draw_height != 0) {
+		if(transparent)
+			draw_transparency_grid();
+
+		device->draw_pixels(rgba, 0, draw_width, draw_height, width, height, transparent);
+	}
 }
 
 bool DisplayBuffer::draw_ready()
diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h
index 1922f875d86..e99fedb0dff 100644
--- a/intern/cycles/render/buffers.h
+++ b/intern/cycles/render/buffers.h
@@ -47,6 +47,7 @@ public:
 	~RenderBuffers();
 
 	void reset(Device *device, int width, int height);
+	float4 *copy_from_device(float exposure, int pass);
 
 protected:
 	void device_free();
@@ -67,6 +68,8 @@ public:
 	   with progressive render we can be using only a subset of the buffer.
 	   if these are zero, it means nothing can be drawn yet */
 	int draw_width, draw_height;
+	/* draw alpha channel? */
+	bool transparent;
 	/* byte buffer for tonemapped result */
 	device_vector<uchar4> rgba;
 	/* mutex, must be locked manually by callers */
@@ -83,6 +86,7 @@ public:
 	bool draw_ready();
 
 protected:
+	void draw_transparency_grid();
 	void device_free();
 
 	Device *device;
diff --git a/intern/cycles/render/camera.h b/intern/cycles/render/camera.h
index 6ae8c8fa4ae..d385aa274f4 100644
--- a/intern/cycles/render/camera.h
+++ b/intern/cycles/render/camera.h
@@ -19,8 +19,6 @@
 #ifndef __CAMERA_H__
 #define __CAMERA_H__
 
-#include "film_response.h"
-
 #include "util_transform.h"
 #include "util_types.h"
 
diff --git a/intern/cycles/render/film.cpp b/intern/cycles/render/film.cpp
index 05874244605..657abd97f70 100644
--- a/intern/cycles/render/film.cpp
+++ b/intern/cycles/render/film.cpp
@@ -19,30 +19,13 @@
 #include "camera.h"
 #include "device.h"
 #include "film.h"
-#include "film_response.h"
 #include "scene.h"
 
 CCL_NAMESPACE_BEGIN
 
-static FilmResponseType find_response_type(const string& name)
-{
-	if(name == "" || name == "None")
-		return FILM_RESPONSE_NONE;
-
-	for(size_t i = 0; i < FILM_RESPONSE_NUM; i++) {
-		FilmResponse *curve = &FILM_RESPONSE[i];
-		if(curve->name == name)
-			return (FilmResponseType)i;
-	}
-
-	return FILM_RESPONSE_NONE;
-}
-
 Film::Film()
 {
 	exposure = 0.8f;
-	response = "Advantix 400";
-	last_response = "";
 	need_update = true;
 }
 
@@ -50,18 +33,6 @@ Film::~Film()
 {
 }
 
-#if 0
-static void generate_python_enum()
-{
-	for(size_t i = 0; i < FILM_RESPONSE_NUM; i++) {
-		FilmResponse *curve = &FILM_RESPONSE[i];
-		/*if(i == 0 || strcmp(curve->brand, FILM_RESPONSE[i-1].brand))
-			printf("(\"\", \"%s\", \"\"),\n", curve->brand);*/
-		printf("(\"%s\", \"%s %s\", \"\"),\n", curve->name, curve->brand, curve->name);
-	}
-}
-#endif
-
 void Film::device_update(Device *device, DeviceScene *dscene)
 {
 	if(!need_update)
@@ -72,57 +43,16 @@ void Film::device_update(Device *device, DeviceScene *dscene)
 	/* update __data */
 	kfilm->exposure = exposure;
 
-	FilmResponseType response_type = find_response_type(response);
-
-	/* update __response_curves */
-	if(response != last_response)  {
-		device_free(device, dscene);
-
-		if(response_type != FILM_RESPONSE_NONE)  {
-			FilmResponse *curve = &FILM_RESPONSE[response_type];
-			size_t response_curve_size = FILM_RESPONSE_SIZE;
-
-			dscene->response_curve_R.copy(curve->B_R, response_curve_size);
-
-			if(curve->rgb) {
-				dscene->response_curve_G.copy(curve->B_G, response_curve_size);
-				dscene->response_curve_B.copy(curve->B_B, response_curve_size);
-			}
-			else {
-				dscene->response_curve_G.copy(curve->B_R, response_curve_size);
-				dscene->response_curve_B.copy(curve->B_R, response_curve_size);
-			}
-
-			device->tex_alloc("__response_curve_R", dscene->response_curve_R, true);
-			device->tex_alloc("__response_curve_G", dscene->response_curve_G, true);
-			device->tex_alloc("__response_curve_B", dscene->response_curve_B, true);
-		}
-
-		last_response = response;
-	}
-
-	kfilm->use_response_curve = (response_type != FILM_RESPONSE_NONE);
-
 	need_update = false;
 }
 
 void Film::device_free(Device *device, DeviceScene *dscene)
 {
-	device->tex_free(dscene->response_curve_R);
-	device->tex_free(dscene->response_curve_G);
-	device->tex_free(dscene->response_curve_B);
-
-	dscene->response_curve_R.clear();
-	dscene->response_curve_G.clear();
-	dscene->response_curve_B.clear();
-
-	last_response = "";
 }
 
 bool Film::modified(const Film& film)
 {
-	return !(response == film.response &&
-		exposure == film.exposure &&
+	return !(exposure == film.exposure &&
 		pass == film.pass);
 }
 
diff --git a/intern/cycles/render/film.h b/intern/cycles/render/film.h
index 9d6c58b30d1..201fc174f4f 100644
--- a/intern/cycles/render/film.h
+++ b/intern/cycles/render/film.h
@@ -29,8 +29,6 @@ class Scene;
 
 class Film {
 public:
-	string response;
-	string last_response;
 	float exposure;
 	int pass;
 	bool need_update;
diff --git a/intern/cycles/render/filter.cpp b/intern/cycles/render/filter.cpp
index d7682bedc12..4925521e4a5 100644
--- a/intern/cycles/render/filter.cpp
+++ b/intern/cycles/render/filter.cpp
@@ -125,5 +125,16 @@ void Filter::device_free(Device *device, DeviceScene *dscene)
 	dscene->filter_table.clear();
 }
 
+bool Filter::modified(const Filter& filter)
+{
+	return !(filter_type == filter.filter_type &&
+		filter_width == filter.filter_width);
+}
+
+void Filter::tag_update(Scene *scene)
+{
+	need_update = true;
+}
+
 CCL_NAMESPACE_END
 
diff --git a/intern/cycles/render/filter.h b/intern/cycles/render/filter.h
index 620e6ebe083..5df7bb8fd14 100644
--- a/intern/cycles/render/filter.h
+++ b/intern/cycles/render/filter.h
@@ -21,6 +21,10 @@
 
 CCL_NAMESPACE_BEGIN
 
+class Device;
+class DeviceScene;
+class Scene;
+
 typedef enum FilterType {
 	FILTER_BOX,
 	FILTER_GAUSSIAN
@@ -38,6 +42,9 @@ public:
 
 	void device_update(Device *device, DeviceScene *dscene);
 	void device_free(Device *device, DeviceScene *dscene);
+
+	bool modified(const Filter& filter);
+	void tag_update(Scene *scene);
 };
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/render/mesh.cpp b/intern/cycles/render/mesh.cpp
index aae32851131..b9163309652 100644
--- a/intern/cycles/render/mesh.cpp
+++ b/intern/cycles/render/mesh.cpp
@@ -392,7 +392,7 @@ void MeshManager::device_update_attributes(Device *device, DeviceScene *dscene,
 	progress.set_status("Updating Mesh", "Computing attributes");
 
 	/* gather per mesh requested attributes. as meshes may have multiple
-	 * shaders assigned, this merged the requested attributes that have
+	 * shaders assigned, this merges the requested attributes that have
 	 * been set per shader by the shader manager */
 	vector<AttributeRequestSet> mesh_attributes(scene->meshes.size());
 
diff --git a/intern/cycles/render/nodes.cpp b/intern/cycles/render/nodes.cpp
index 546d74c5c22..5ffc3fbbabb 100644
--- a/intern/cycles/render/nodes.cpp
+++ b/intern/cycles/render/nodes.cpp
@@ -1275,6 +1275,24 @@ void BackgroundNode::compile(OSLCompiler& compiler)
 	compiler.add(this, "node_background");
 }
 
+/* Holdout Closure */
+
+HoldoutNode::HoldoutNode()
+: ShaderNode("holdout")
+{
+	add_output("Holdout", SHADER_SOCKET_CLOSURE);
+}
+
+void HoldoutNode::compile(SVMCompiler& compiler)
+{
+	compiler.add_node(NODE_CLOSURE_HOLDOUT, CLOSURE_HOLDOUT_ID);
+}
+
+void HoldoutNode::compile(OSLCompiler& compiler)
+{
+	compiler.add(this, "node_holdout");
+}
+
 /* Geometry */
 
 GeometryNode::GeometryNode()
diff --git a/intern/cycles/render/nodes.h b/intern/cycles/render/nodes.h
index da3b743e79b..2afe585c3ac 100644
--- a/intern/cycles/render/nodes.h
+++ b/intern/cycles/render/nodes.h
@@ -257,6 +257,11 @@ public:
 	SHADER_NODE_CLASS(BackgroundNode)
 };
 
+class HoldoutNode : public ShaderNode {
+public:
+	SHADER_NODE_CLASS(HoldoutNode)
+};
+
 class GeometryNode : public ShaderNode {
 public:
 	SHADER_NODE_CLASS(GeometryNode)
diff --git a/intern/cycles/render/scene.h b/intern/cycles/render/scene.h
index 965a0054846..9d475cd9fa5 100644
--- a/intern/cycles/render/scene.h
+++ b/intern/cycles/render/scene.h
@@ -84,11 +84,6 @@ public:
 	/* filter */
 	device_vector<float> filter_table;
 
-	/* film */
-	device_vector<float> response_curve_R;
-	device_vector<float> response_curve_G;
-	device_vector<float> response_curve_B;
-
 	/* integrator */
 	device_vector<uint> sobol_directions;
 
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index 56fbadcee08..5890714d740 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -52,6 +52,7 @@ Session::Session(const SessionParams& params_)
 	delayed_reset.do_reset = false;
 	delayed_reset.w = 0;
 	delayed_reset.h = 0;
+	delayed_reset.passes = 0;
 
 	display_outdated = false;
 	gpu_draw_ready = false;
@@ -95,7 +96,7 @@ bool Session::ready_to_reset()
 
 /* GPU Session */
 
-void Session::reset_gpu(int w, int h)
+void Session::reset_gpu(int w, int h, int passes)
 {
 	/* block for buffer acces and reset immediately. we can't do this
 	   in the thread, because we need to allocate an OpenGL buffer, and
@@ -106,7 +107,7 @@ void Session::reset_gpu(int w, int h)
 	display_outdated = true;
 	reset_time = time_dt();
 
-	reset_(w, h);
+	reset_(w, h, passes);
 
 	gpu_need_tonemap = false;
 	gpu_need_tonemap_cond.notify_all();
@@ -148,7 +149,14 @@ void Session::run_gpu()
 	start_time = time_dt();
 	reset_time = time_dt();
 
-	while(!progress.get_cancel() && tile_manager.next()) {
+	while(!progress.get_cancel()) {
+		bool done = !tile_manager.next();
+
+		if(done && params.background)
+			break;
+
+		/* todo: wait when done in interactive mode */
+
 		/* buffers mutex is locked entirely while rendering each
 		   pass, and released/reacquired on each iteration to allow
 		   reset and draw in between */
@@ -193,7 +201,7 @@ void Session::run_gpu()
 
 /* CPU Session */
 
-void Session::reset_cpu(int w, int h)
+void Session::reset_cpu(int w, int h, int passes)
 {
 	thread_scoped_lock reset_lock(delayed_reset.mutex);
 
@@ -202,6 +210,7 @@ void Session::reset_cpu(int w, int h)
 
 	delayed_reset.w = w;
 	delayed_reset.h = h;
+	delayed_reset.passes = passes;
 	delayed_reset.do_reset = true;
 	device->task_cancel();
 }
@@ -235,11 +244,18 @@ void Session::run_cpu()
 		thread_scoped_lock buffers_lock(buffers->mutex);
 		thread_scoped_lock display_lock(display->mutex);
 
-		reset_(delayed_reset.w, delayed_reset.h);
+		reset_(delayed_reset.w, delayed_reset.h, delayed_reset.passes);
 		delayed_reset.do_reset = false;
 	}
 
-	while(!progress.get_cancel() && tile_manager.next()) {
+	while(!progress.get_cancel()) {
+		bool done = !tile_manager.next();
+
+		if(done && params.background)
+			break;
+
+		/* todo: wait when done in interactive mode */
+
 		{
 			thread_scoped_lock buffers_lock(buffers->mutex);
 
@@ -266,7 +282,7 @@ void Session::run_cpu()
 			if(delayed_reset.do_reset) {
 				/* reset rendering if request from main thread */
 				delayed_reset.do_reset = false;
-				reset_(delayed_reset.w, delayed_reset.h);
+				reset_(delayed_reset.w, delayed_reset.h, delayed_reset.passes);
 			}
 			else {
 				/* tonemap only if we do not reset, we don't we don't
@@ -313,7 +329,7 @@ bool Session::draw(int w, int h)
 		return draw_cpu(w, h);
 }
 
-void Session::reset_(int w, int h)
+void Session::reset_(int w, int h, int passes)
 {
 	if(w != buffers->width || h != buffers->height) {
 		gpu_draw_ready = false;
@@ -321,19 +337,28 @@ void Session::reset_(int w, int h)
 		display->reset(device, w, h);
 	}
 
-	tile_manager.reset(w, h);
+	tile_manager.reset(w, h, passes);
 
 	start_time = time_dt();
 	preview_time = 0.0;
 	pass = 0;
 }
 
-void Session::reset(int w, int h)
+void Session::reset(int w, int h, int passes)
 {
 	if(device_use_gl)
-		reset_gpu(w, h);
+		reset_gpu(w, h, passes);
 	else
-		reset_cpu(w, h);
+		reset_cpu(w, h, passes);
+}
+
+void Session::set_passes(int passes)
+{
+	if(passes != params.passes) {
+		params.passes = passes;
+		tile_manager.set_passes(passes);
+		/* todo: awake in paused loop */
+	}
 }
 
 void Session::wait()
diff --git a/intern/cycles/render/session.h b/intern/cycles/render/session.h
index 3411dcd23c6..18983f55353 100644
--- a/intern/cycles/render/session.h
+++ b/intern/cycles/render/session.h
@@ -107,13 +107,15 @@ public:
 	void wait();
 
 	bool ready_to_reset();
-	void reset(int w, int h);
+	void reset(int w, int h, int passes);
+	void set_passes(int passes);
 
 protected:
 	struct DelayedReset {
 		thread_mutex mutex;
 		bool do_reset;
 		int w, h;
+		int passes;
 	} delayed_reset;
 
 	void run();
@@ -123,15 +125,15 @@ protected:
 
 	void tonemap();
 	void path_trace(Tile& tile);
-	void reset_(int w, int h);
+	void reset_(int w, int h, int passes);
 
 	void run_cpu();
 	bool draw_cpu(int w, int h);
-	void reset_cpu(int w, int h);
+	void reset_cpu(int w, int h, int passes);
 
 	void run_gpu();
 	bool draw_gpu(int w, int h);
-	void reset_gpu(int w, int h);
+	void reset_gpu(int w, int h, int passes);
 
 	TileManager tile_manager;
 	bool device_use_gl;
diff --git a/intern/cycles/render/tile.cpp b/intern/cycles/render/tile.cpp
index 28ea39e70ed..450090c42f8 100644
--- a/intern/cycles/render/tile.cpp
+++ b/intern/cycles/render/tile.cpp
@@ -25,18 +25,17 @@ CCL_NAMESPACE_BEGIN
 TileManager::TileManager(bool progressive_, int passes_, int tile_size_, int min_size_)
 {
 	progressive = progressive_;
-	passes = passes_;
 	tile_size = tile_size_;
 	min_size = min_size_;
 
-	reset(0, 0);
+	reset(0, 0, 0);
 }
 
 TileManager::~TileManager()
 {
 }
 
-void TileManager::reset(int width_, int height_)
+void TileManager::reset(int width_, int height_, int passes_)
 {
 	full_width = width_;
 	full_height = height_;
@@ -54,6 +53,8 @@ void TileManager::reset(int width_, int height_)
 		}
 	}
 
+	passes = passes_;
+
 	state.width = 0;
 	state.height = 0;
 	state.pass = -1;
@@ -61,6 +62,11 @@ void TileManager::reset(int width_, int height_)
 	state.tiles.clear();
 }
 
+void TileManager::set_passes(int passes_)
+{
+	passes = passes_;
+}
+
 void TileManager::set_tiles()
 {
 	int resolution = state.resolution;
diff --git a/intern/cycles/render/tile.h b/intern/cycles/render/tile.h
index c3517b1e24d..56c69cdce88 100644
--- a/intern/cycles/render/tile.h
+++ b/intern/cycles/render/tile.h
@@ -50,7 +50,8 @@ public:
 	TileManager(bool progressive, int passes, int tile_size, int min_size);
 	~TileManager();
 
-	void reset(int width, int height);
+	void reset(int width, int height, int passes);
+	void set_passes(int passes);
 	bool next();
 	bool done();
 
diff --git a/intern/cycles/util/util_color.h b/intern/cycles/util/util_color.h
index a11a5c7c2a7..8b4a175f498 100644
--- a/intern/cycles/util/util_color.h
+++ b/intern/cycles/util/util_color.h
@@ -60,6 +60,11 @@ __device float3 color_scene_linear_to_srgb(float3 c)
 
 #endif
 
+__device float linear_rgb_to_gray(float3 c)
+{
+	return c.x*0.2126f + c.y*0.7152f + c.z*0.0722f;
+}
+
 CCL_NAMESPACE_END
 
 #endif /* __UTIL_COLOR_H__ */
diff --git a/intern/cycles/util/util_types.h b/intern/cycles/util/util_types.h
index 12ae60f1143..af5b91c963a 100644
--- a/intern/cycles/util/util_types.h
+++ b/intern/cycles/util/util_types.h
@@ -30,6 +30,7 @@
 #ifndef __KERNEL_GPU__
 
 #define __device static inline
+#define __device_noinline static
 #define __global
 #define __local
 #define __shared
diff --git a/source/blender/blenkernel/BKE_node.h b/source/blender/blenkernel/BKE_node.h
index f61231141fd..5b1a80186a6 100644
--- a/source/blender/blenkernel/BKE_node.h
+++ b/source/blender/blenkernel/BKE_node.h
@@ -311,6 +311,7 @@ struct ShadeResult;
 #define SH_NODE_ADD_CLOSURE				156
 #define SH_NODE_TEX_ENVIRONMENT			157
 #define SH_NODE_OUTPUT_TEXTURE			158
+#define SH_NODE_HOLDOUT					159
 
 /* custom defines options for Material node */
 #define SH_NODE_MAT_DIFF   1
diff --git a/source/blender/blenkernel/intern/node.c b/source/blender/blenkernel/intern/node.c
index 927d406eb61..a36032f137e 100644
--- a/source/blender/blenkernel/intern/node.c
+++ b/source/blender/blenkernel/intern/node.c
@@ -3702,10 +3702,11 @@ static void registerShaderNodes(ListBase *ntypelist)
 	register_node_type_sh_emission(ntypelist);
 	register_node_type_sh_mix_closure(ntypelist);
 	register_node_type_sh_add_closure(ntypelist);
+	register_node_type_sh_holdout(ntypelist);
 
 	register_node_type_sh_output_lamp(ntypelist);
 	register_node_type_sh_output_material(ntypelist);
-	register_node_type_sh_output_texture(ntypelist);
+	//register_node_type_sh_output_texture(ntypelist);
 	register_node_type_sh_output_world(ntypelist);
 
 	register_node_type_sh_tex_blend(ntypelist);
diff --git a/source/blender/makesrna/intern/rna_nodetree_types.h b/source/blender/makesrna/intern/rna_nodetree_types.h
index 41f1bf2d383..d14e4cd76f2 100644
--- a/source/blender/makesrna/intern/rna_nodetree_types.h
+++ b/source/blender/makesrna/intern/rna_nodetree_types.h
@@ -52,7 +52,7 @@ DefNode( ShaderNode,     SH_NODE_COMBRGB,         0,                      "COMBR
 DefNode( ShaderNode,     SH_NODE_HUE_SAT,         0,                      "HUE_SAT",        HueSaturation,    "Hue/Saturation",    ""              )
 DefNode( ShaderNode,     SH_NODE_OUTPUT_MATERIAL, 0,                      "OUTPUT_MATERIAL",OutputMaterial,   "Material Output",   ""              )
 DefNode( ShaderNode,     SH_NODE_OUTPUT_LAMP,     0,                      "OUTPUT_LAMP",    OutputLamp,       "Lamp Output",       ""              )
-DefNode( ShaderNode,     SH_NODE_OUTPUT_TEXTURE,  0,                      "OUTPUT_TEXTURE", OutputTexture,    "Texture Output",    ""              )
+//DefNode( ShaderNode,     SH_NODE_OUTPUT_TEXTURE,  0,                      "OUTPUT_TEXTURE", OutputTexture,    "Texture Output",    ""              )
 DefNode( ShaderNode,     SH_NODE_OUTPUT_WORLD,    0,                      "OUTPUT_WORLD",   OutputWorld,      "World Output",      ""              )
 DefNode( ShaderNode,     SH_NODE_FRESNEL,         0,                      "FRESNEL",        Fresnel,          "Fresnel",           ""              )
 DefNode( ShaderNode,     SH_NODE_MIX_CLOSURE,     0,                      "MIX_CLOSURE",    MixClosure,       "Mix Closure",       ""              )
@@ -60,6 +60,7 @@ DefNode( ShaderNode,     SH_NODE_ADD_CLOSURE,     0,                      "ADD_C
 
 DefNode( ShaderNode,     SH_NODE_ATTRIBUTE, 0, "ATTRIBUTE", Attribute, "Attribute", "")
 DefNode( ShaderNode,     SH_NODE_BACKGROUND, 0, "BACKGROUND", Background, "Background", "")
+DefNode( ShaderNode,     SH_NODE_HOLDOUT, 0, "HOLDOUT", Holdout, "Holdout", "")
 DefNode( ShaderNode,     SH_NODE_BSDF_ANISOTROPIC, 0, "BSDF_ANISOTROPIC", BsdfAnisotropic, "Bsdf Anisotropic", "")
 DefNode( ShaderNode,     SH_NODE_BSDF_DIFFUSE, 0, "BSDF_DIFFUSE", BsdfDiffuse, "Diffuse Bsdf", "")
 DefNode( ShaderNode,     SH_NODE_BSDF_GLOSSY, def_glossy, "BSDF_GLOSSY", BsdfGlossy, "Glossy Bsdf", "")
diff --git a/source/blender/nodes/CMakeLists.txt b/source/blender/nodes/CMakeLists.txt
index 96c59992b25..f4b6a452281 100644
--- a/source/blender/nodes/CMakeLists.txt
+++ b/source/blender/nodes/CMakeLists.txt
@@ -128,6 +128,7 @@ set(SRC
 	intern/SHD_nodes/SHD_emission.c
 	intern/SHD_nodes/SHD_fresnel.c
 	intern/SHD_nodes/SHD_geometry.c
+	intern/SHD_nodes/SHD_holdout.c
 	intern/SHD_nodes/SHD_light_path.c
 	intern/SHD_nodes/SHD_mix_closure.c
 	intern/SHD_nodes/SHD_add_closure.c
diff --git a/source/blender/nodes/SHD_node.h b/source/blender/nodes/SHD_node.h
index 7bee0dfb0bb..a939bca161d 100644
--- a/source/blender/nodes/SHD_node.h
+++ b/source/blender/nodes/SHD_node.h
@@ -67,6 +67,7 @@ void register_node_type_sh_bsdf_translucent(ListBase *lb);
 void register_node_type_sh_bsdf_transparent(ListBase *lb);
 void register_node_type_sh_bsdf_velvet(ListBase *lb);
 void register_node_type_sh_emission(ListBase *lb);
+void register_node_type_sh_holdout(ListBase *lb);
 void register_node_type_sh_mix_closure(ListBase *lb);
 void register_node_type_sh_add_closure(ListBase *lb);
 
diff --git a/source/blender/nodes/intern/SHD_nodes/SHD_holdout.c b/source/blender/nodes/intern/SHD_nodes/SHD_holdout.c
new file mode 100644
index 00000000000..a96b621456a
--- /dev/null
+++ b/source/blender/nodes/intern/SHD_nodes/SHD_holdout.c
@@ -0,0 +1,63 @@
+/**
+ * $Id: SHD_output.c 32517 2010-10-16 14:32:17Z campbellbarton $
+ *
+ * ***** BEGIN GPL LICENSE BLOCK *****
+ *
+ * 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.
+ *
+ * The Original Code is Copyright (C) 2005 Blender Foundation.
+ * All rights reserved.
+ *
+ * The Original Code is: all of this file.
+ *
+ * Contributor(s): none yet.
+ *
+ * ***** END GPL LICENSE BLOCK *****
+ */
+
+#include "../SHD_util.h"
+
+/* **************** OUTPUT ******************** */
+
+static bNodeSocketType sh_node_holdout_in[]= {
+	{	-1, 0, ""	}
+};
+
+static bNodeSocketType sh_node_holdout_out[]= {
+	{	SOCK_CLOSURE, 0, "Holdout",		0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f},
+	{	-1, 0, ""	}
+};
+
+static void node_shader_exec_holdout(void *UNUSED(data), bNode *UNUSED(node), bNodeStack **UNUSED(in), bNodeStack **UNUSED(out))
+{
+}
+
+
+/* node type definition */
+void register_node_type_sh_holdout(ListBase *lb)
+{
+	static bNodeType ntype;
+
+	node_type_base(&ntype, SH_NODE_HOLDOUT, "Holdout", NODE_CLASS_CLOSURE, 0,
+		sh_node_holdout_in, sh_node_holdout_out);
+	node_type_size(&ntype, 150, 60, 200);
+	node_type_init(&ntype, NULL);
+	node_type_storage(&ntype, "", NULL, NULL);
+	node_type_exec(&ntype, node_shader_exec_holdout);
+	node_type_gpu(&ntype, NULL);
+
+	nodeRegisterType(lb, &ntype);
+};
+
-- 
cgit v1.2.3