Merge branch 'master' of git.blender.org:blender into cycles_embree

# Conflicts:
#	build_files/cmake/platform/platform_win32.cmake

73 files changed, 1362 insertions, 318 deletions

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 79309ac0f7b..817588c12fe 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -578,6 +578,9 @@ if(WIN32)
 
 	option(WINDOWS_USE_VISUAL_STUDIO_FOLDERS "Organize the visual studio project according to source folders." ON)
 	mark_as_advanced(WINDOWS_USE_VISUAL_STUDIO_FOLDERS)
+
+	option(WINDOWS_PYTHON_DEBUG "Include the files needed for debugging python scripts with visual studio 2017+." OFF)
+	mark_as_advanced(WINDOWS_PYTHON_DEBUG)
 endif()
 
 # avoid using again
diff --git a/build_files/build_environment/cmake/versions.cmake b/build_files/build_environment/cmake/versions.cmake
index 7ede80a2be9..7bd994c3491 100644
--- a/build_files/build_environment/cmake/versions.cmake
+++ b/build_files/build_environment/cmake/versions.cmake
@@ -21,7 +21,7 @@ set(ZLIB_URI https://zlib.net/zlib-${ZLIB_VERSION}.tar.gz)
 set(ZLIB_HASH 1c9f62f0778697a09d36121ead88e08e)
 
 set(OPENAL_VERSION 1.18.2)
-set(OPENAL_URI http://kcat.strangesoft.net/openal-releases/openal-soft-${OPENAL_VERSION}.tar.bz2)
+set(OPENAL_URI http://openal-soft.org/openal-releases/openal-soft-${OPENAL_VERSION}.tar.bz2)
 set(OPENAL_HASH d4eeb0889812e2fdeaa1843523d76190)
 
 set(PNG_VERSION 1.6.35)
diff --git a/build_files/cmake/macros.cmake b/build_files/cmake/macros.cmake
index 677233190cf..2b4d454a6a7 100644
--- a/build_files/cmake/macros.cmake
+++ b/build_files/cmake/macros.cmake
@@ -669,10 +669,10 @@ function(SETUP_BLENDER_SORTED_LIBS)
 		bf_intern_mikktspace
 		bf_intern_dualcon
 		bf_intern_cycles
+		cycles_device
 		cycles_render
 		cycles_graph
 		cycles_bvh
-		cycles_device
 		cycles_kernel
 		cycles_util
 		cycles_subd
@@ -1198,7 +1198,11 @@ function(delayed_do_install
 		foreach(i RANGE ${n})
 			list(GET files ${i} f)
 			list(GET destinations ${i} d)
-			install(FILES ${f} DESTINATION ${targetdir}/${d})
+			if(NOT IS_ABSOLUTE ${d})
+				install(FILES ${f} DESTINATION ${targetdir}/${d})
+			else()
+				install(FILES ${f} DESTINATION ${d})
+			endif()
 		endforeach()
 	endif()
 endfunction()
@@ -1235,6 +1239,8 @@ function(data_to_c_simple
 	get_filename_component(_file_to   ${CMAKE_CURRENT_BINARY_DIR}/${file_from}.c REALPATH)
 
 	list(APPEND ${list_to_add} ${_file_to})
+	source_group(Generated FILES ${_file_to})
+	list(APPEND ${list_to_add} ${file_from})
 	set(${list_to_add} ${${list_to_add}} PARENT_SCOPE)
 
 	get_filename_component(_file_to_path ${_file_to} PATH)
diff --git a/build_files/cmake/platform/platform_win32.cmake b/build_files/cmake/platform/platform_win32.cmake
index bede5de4a71..38d3db02a32 100644
--- a/build_files/cmake/platform/platform_win32.cmake
+++ b/build_files/cmake/platform/platform_win32.cmake
@@ -611,3 +611,40 @@ endif()
 if(WITH_CYCLES_EMBREE)
 	find_package(embree 3.2 REQUIRED)
 endif()
+
+if (WINDOWS_PYTHON_DEBUG)
+	# Include the system scripts in the blender_python_system_scripts project.
+	FILE(GLOB_RECURSE inFiles "${CMAKE_SOURCE_DIR}/release/scripts/*.*" )
+	ADD_CUSTOM_TARGET(blender_python_system_scripts SOURCES ${inFiles})
+	foreach(_source IN ITEMS ${inFiles})
+		get_filename_component(_source_path "${_source}" PATH)
+		string(REPLACE "${CMAKE_SOURCE_DIR}/release/scripts/" "" _source_path "${_source_path}")
+		string(REPLACE "/" "\\" _group_path "${_source_path}")
+		source_group("${_group_path}" FILES "${_source}")
+	endforeach()
+	# Include the user scripts from the profile folder in the blender_python_user_scripts project.
+	set(USER_SCRIPTS_ROOT "$ENV{appdata}/blender foundation/blender/${BLENDER_VERSION}")
+	file(TO_CMAKE_PATH ${USER_SCRIPTS_ROOT} USER_SCRIPTS_ROOT)
+	FILE(GLOB_RECURSE inFiles "${USER_SCRIPTS_ROOT}/scripts/*.*" )
+	ADD_CUSTOM_TARGET(blender_python_user_scripts SOURCES ${inFiles})
+	foreach(_source IN ITEMS ${inFiles})
+		get_filename_component(_source_path "${_source}" PATH)
+		string(REPLACE "${USER_SCRIPTS_ROOT}/scripts" "" _source_path "${_source_path}")
+		string(REPLACE "/" "\\" _group_path "${_source_path}")
+		source_group("${_group_path}" FILES "${_source}")
+	endforeach()
+	set_target_properties(blender_python_system_scripts PROPERTIES FOLDER "scripts")
+	set_target_properties(blender_python_user_scripts PROPERTIES FOLDER "scripts")
+	# Set the default debugging options for the project, only write this file once so the user
+	# is free to override them at their own perril.
+	set(USER_PROPS_FILE "${CMAKE_CURRENT_BINARY_DIR}/source/creator/blender.Cpp.user.props")
+	if(NOT EXISTS ${USER_PROPS_FILE})
+		# Layout below is messy, because otherwise the generated file will look messy.
+		file(WRITE ${USER_PROPS_FILE} "<?xml version=\"1.0\" encoding=\"utf-8\"?>
+<Project DefaultTargets=\"Build\" xmlns=\"http://schemas.microsoft.com/developer/msbuild/2003\">
+	<PropertyGroup>
+		<LocalDebuggerCommandArguments>-con --env-system-scripts \"${CMAKE_SOURCE_DIR}/release/scripts\" </LocalDebuggerCommandArguments>
+	</PropertyGroup>
+</Project>")
+	endif()
+endif()
diff --git a/build_files/windows/configure_msbuild.cmd b/build_files/windows/configure_msbuild.cmd
index f8c2a87de8e..a0ffe0c58fd 100644
--- a/build_files/windows/configure_msbuild.cmd
+++ b/build_files/windows/configure_msbuild.cmd
@@ -19,7 +19,11 @@ if "%WITH_CLANG%"=="1" (
 		exit /b 1 
 	)
 )
-set BUILD_CMAKE_ARGS=%BUILD_CMAKE_ARGS% -G "Visual Studio %BUILD_VS_VER% %BUILD_VS_YEAR%%WINDOWS_ARCH%" %TESTS_CMAKE_ARGS% %CLANG_CMAKE_ARGS% %ASAN_CMAKE_ARGS%
+
+if "%WITH_PYDEBUG%"=="1" (
+	set PYDEBUG_CMAKE_ARGS=-DWINDOWS_PYTHON_DEBUG=On
+)
+set BUILD_CMAKE_ARGS=%BUILD_CMAKE_ARGS% -G "Visual Studio %BUILD_VS_VER% %BUILD_VS_YEAR%%WINDOWS_ARCH%" %TESTS_CMAKE_ARGS% %CLANG_CMAKE_ARGS% %ASAN_CMAKE_ARGS% %PYDEBUG_CMAKE_ARGS%
 
 if NOT EXIST %BUILD_DIR%\nul (
 	mkdir %BUILD_DIR%
diff --git a/build_files/windows/parse_arguments.cmd b/build_files/windows/parse_arguments.cmd
index e8caddaf4ea..30bd2ebdcd2 100644
--- a/build_files/windows/parse_arguments.cmd
+++ b/build_files/windows/parse_arguments.cmd
@@ -64,6 +64,8 @@ if NOT "%1" == "" (
 		shift /1
 	) else if "%1" == "nobuild" (
 		set NOBUILD=1
+	) else if "%1" == "pydebug" (
+		set WITH_PYDEBUG=1
 	) else if "%1" == "showhash" (
 		SET BUILD_SHOW_HASHES=1
 	REM Non-Build Commands
diff --git a/build_files/windows/reset_variables.cmd b/build_files/windows/reset_variables.cmd
index a522ed7407f..ba584a7f34a 100644
--- a/build_files/windows/reset_variables.cmd
+++ b/build_files/windows/reset_variables.cmd
@@ -25,3 +25,5 @@ set WITH_CLANG=
 set WITH_ASAN=
 set CLANG_CMAKE_ARGS=
 set ASAN_CMAKE_ARGS=
+set WITH_PYDEBUG=
+set PYDEBUG_CMAKE_ARGS=
diff --git a/intern/cycles/CMakeLists.txt b/intern/cycles/CMakeLists.txt
index fd330c7580c..6b58ec78d1c 100644
--- a/intern/cycles/CMakeLists.txt
+++ b/intern/cycles/CMakeLists.txt
@@ -3,7 +3,14 @@ if(NOT WITH_BLENDER AND WITH_CYCLES_STANDALONE)
 	set(CYCLES_INSTALL_PATH "")
 else()
 	set(WITH_CYCLES_BLENDER ON)
-	set(CYCLES_INSTALL_PATH "scripts/addons/cycles")
+	# WINDOWS_PYTHON_DEBUG needs to write into the user addons folder since it will
+	# be started with --env-system-scripts pointing to the release folder, which will
+	# lack the cycles addon, and we don't want to write into it.
+	if(NOT WINDOWS_PYTHON_DEBUG)
+		set(CYCLES_INSTALL_PATH "scripts/addons/cycles")
+	else()
+		set(CYCLES_INSTALL_PATH "$ENV{appdata}/blender foundation/blender/${BLENDER_VERSION}/scripts/addons/cycles")
+	endif()
 endif()
 
 # External Libraries
diff --git a/intern/cycles/blender/addon/engine.py b/intern/cycles/blender/addon/engine.py
index 87dcbe486c7..55638415e66 100644
--- a/intern/cycles/blender/addon/engine.py
+++ b/intern/cycles/blender/addon/engine.py
@@ -254,21 +254,32 @@ def register_passes(engine, scene, srl):
     if crl.use_pass_volume_indirect:           engine.register_pass(scene, srl, "VolumeInd",                     3, "RGB", 'COLOR')
 
     cscene = scene.cycles
-    if crl.use_denoising and crl.denoising_store_passes and not cscene.use_progressive_refine:
-        engine.register_pass(scene, srl, "Denoising Normal",          3, "XYZ", 'VECTOR')
-        engine.register_pass(scene, srl, "Denoising Normal Variance", 3, "XYZ", 'VECTOR')
-        engine.register_pass(scene, srl, "Denoising Albedo",          3, "RGB", 'COLOR')
-        engine.register_pass(scene, srl, "Denoising Albedo Variance", 3, "RGB", 'COLOR')
-        engine.register_pass(scene, srl, "Denoising Depth",           1, "Z",   'VALUE')
-        engine.register_pass(scene, srl, "Denoising Depth Variance",  1, "Z",   'VALUE')
-        engine.register_pass(scene, srl, "Denoising Shadow A",        3, "XYV", 'VECTOR')
-        engine.register_pass(scene, srl, "Denoising Shadow B",        3, "XYV", 'VECTOR')
-        engine.register_pass(scene, srl, "Denoising Image",           3, "RGB", 'COLOR')
-        engine.register_pass(scene, srl, "Denoising Image Variance",  3, "RGB", 'COLOR')
-
-        clean_options = ("denoising_diffuse_direct", "denoising_diffuse_indirect",
-                         "denoising_glossy_direct", "denoising_glossy_indirect",
-                         "denoising_transmission_direct", "denoising_transmission_indirect",
-                         "denoising_subsurface_direct", "denoising_subsurface_indirect")
-        if any(getattr(crl, option) for option in clean_options):
-            engine.register_pass(scene, srl, "Denoising Clean", 3, "RGB", 'COLOR')
+
+    if crl.use_pass_crypto_object:
+        for i in range(0, crl.pass_crypto_depth, 2):
+            engine.register_pass(scene, srl, "CryptoObject" + '{:02d}'.format(i), 4, "RGBA", 'COLOR')
+    if crl.use_pass_crypto_material:
+        for i in range(0, crl.pass_crypto_depth, 2):
+            engine.register_pass(scene, srl, "CryptoMaterial" + '{:02d}'.format(i), 4, "RGBA", 'COLOR')
+    if srl.cycles.use_pass_crypto_asset:
+        for i in range(0, srl.cycles.pass_crypto_depth, 2):
+            engine.register_pass(scene, srl, "CryptoAsset" + '{:02d}'.format(i), 4, "RGBA", 'COLOR')
+
+    if crl.use_denoising or crl.denoising_store_passes:
+        engine.register_pass(scene, srl, "Noisy Image", 3, "RGBA", 'COLOR')
+        if crl.denoising_store_passes:
+            engine.register_pass(scene, srl, "Denoising Normal",          3, "XYZ", 'VECTOR')
+            engine.register_pass(scene, srl, "Denoising Normal Variance", 3, "XYZ", 'VECTOR')
+            engine.register_pass(scene, srl, "Denoising Albedo",          3, "RGB", 'COLOR')
+            engine.register_pass(scene, srl, "Denoising Albedo Variance", 3, "RGB", 'COLOR')
+            engine.register_pass(scene, srl, "Denoising Depth",           1, "Z",   'VALUE')
+            engine.register_pass(scene, srl, "Denoising Depth Variance",  1, "Z",   'VALUE')
+            engine.register_pass(scene, srl, "Denoising Shadow A",        3, "XYV", 'VECTOR')
+            engine.register_pass(scene, srl, "Denoising Shadow B",        3, "XYV", 'VECTOR')
+            engine.register_pass(scene, srl, "Denoising Image Variance",  3, "RGB", 'COLOR')
+            clean_options = ("denoising_diffuse_direct", "denoising_diffuse_indirect",
+                             "denoising_glossy_direct", "denoising_glossy_indirect",
+                             "denoising_transmission_direct", "denoising_transmission_indirect",
+                             "denoising_subsurface_direct", "denoising_subsurface_indirect")
+            if any(getattr(crl, option) for option in clean_options):
+                engine.register_pass(scene, srl, "Denoising Clean", 3, "RGB", 'COLOR')
diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index fce2fca15c5..d25a77a52e7 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -1344,7 +1344,36 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
             default=False,
             update=update_render_passes,
         )
-
+        cls.use_pass_crypto_object = BoolProperty(
+                name="Cryptomatte Object",
+                description="Render cryptomatte object pass, for isolating objects in compositing",
+                default=False,
+                update=update_render_passes,
+                )
+        cls.use_pass_crypto_material = BoolProperty(
+                name="Cryptomatte Material",
+                description="Render cryptomatte material pass, for isolating materials in compositing",
+                default=False,
+                update=update_render_passes,
+                )
+        cls.use_pass_crypto_asset = BoolProperty(
+                name="Cryptomatte Asset",
+                description="Render cryptomatte asset pass, for isolating groups of objects with the same parent",
+                default=False,
+                update=update_render_passes,
+                )
+        cls.pass_crypto_depth = IntProperty(
+                name="Cryptomatte Levels",
+                description="Sets how many unique objects can be distinguished per pixel",
+                default=6, min=2, max=16, step=2,
+                update=update_render_passes,
+                )
+        cls.pass_crypto_accurate = BoolProperty(
+                name="Cryptomatte Accurate",
+                description="Gerenate a more accurate Cryptomatte pass. CPU only, may render slower and use more memory",
+                default=True,
+                update=update_render_passes,
+                )
     @classmethod
     def unregister(cls):
         del bpy.types.SceneRenderLayer.cycles
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 006554fdb45..017c5735176 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -523,6 +523,8 @@ class CYCLES_RENDER_PT_layer_passes(CyclesButtonsPanel, Panel):
         col.prop(rl, "use_pass_shadow")
         col.prop(rl, "use_pass_ambient_occlusion")
         col.separator()
+        col.prop(crl, "denoising_store_passes", text="Denoising Data")
+        col.separator()
         col.prop(rl, "pass_alpha_threshold")
 
         col = split.column()
@@ -555,12 +557,6 @@ class CYCLES_RENDER_PT_layer_passes(CyclesButtonsPanel, Panel):
         col.prop(rl, "use_pass_emit", text="Emission")
         col.prop(rl, "use_pass_environment")
 
-        if context.scene.cycles.feature_set == 'EXPERIMENTAL':
-            col.separator()
-            sub = col.column()
-            sub.active = crl.use_denoising
-            sub.prop(crl, "denoising_store_passes", text="Denoising")
-
         col = layout.column()
         col.prop(crl, "pass_debug_render_time")
         if _cycles.with_cycles_debug:
@@ -569,6 +565,17 @@ class CYCLES_RENDER_PT_layer_passes(CyclesButtonsPanel, Panel):
             col.prop(crl, "pass_debug_bvh_intersections")
             col.prop(crl, "pass_debug_ray_bounces")
 
+        crl = rl.cycles
+        layout.label("Cryptomatte:")
+        row = layout.row(align=True)
+        row.prop(crl, "use_pass_crypto_object", text="Object", toggle=True)
+        row.prop(crl, "use_pass_crypto_material", text="Material", toggle=True)
+        row.prop(crl, "use_pass_crypto_asset", text="Asset", toggle=True)
+        row = layout.row(align=True)
+        row.prop(crl, "pass_crypto_depth")
+        row = layout.row(align=True)
+        row.active = use_cpu(context)
+        row.prop(crl, "pass_crypto_accurate", text="Accurate Mode")
 
 class CYCLES_RENDER_PT_views(CyclesButtonsPanel, Panel):
     bl_label = "Views"
@@ -636,9 +643,8 @@ class CYCLES_RENDER_PT_denoising(CyclesButtonsPanel, Panel):
         rl = rd.layers.active
         crl = rl.cycles
 
-        layout.active = crl.use_denoising
-
         split = layout.split()
+        split.active = crl.use_denoising
 
         col = split.column()
         sub = col.column(align=True)
@@ -653,24 +659,28 @@ class CYCLES_RENDER_PT_denoising(CyclesButtonsPanel, Panel):
         layout.separator()
 
         row = layout.row()
+        row.active = crl.use_denoising or crl.denoising_store_passes
         row.label(text="Diffuse:")
         sub = row.row(align=True)
         sub.prop(crl, "denoising_diffuse_direct", text="Direct", toggle=True)
         sub.prop(crl, "denoising_diffuse_indirect", text="Indirect", toggle=True)
 
         row = layout.row()
+        row.active = crl.use_denoising or crl.denoising_store_passes
         row.label(text="Glossy:")
         sub = row.row(align=True)
         sub.prop(crl, "denoising_glossy_direct", text="Direct", toggle=True)
         sub.prop(crl, "denoising_glossy_indirect", text="Indirect", toggle=True)
 
         row = layout.row()
+        row.active = crl.use_denoising or crl.denoising_store_passes
         row.label(text="Transmission:")
         sub = row.row(align=True)
         sub.prop(crl, "denoising_transmission_direct", text="Direct", toggle=True)
         sub.prop(crl, "denoising_transmission_indirect", text="Indirect", toggle=True)
 
         row = layout.row()
+        row.active = crl.use_denoising or crl.denoising_store_passes
         row.label(text="Subsurface:")
         sub = row.row(align=True)
         sub.prop(crl, "denoising_subsurface_direct", text="Direct", toggle=True)
diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp
index 0fab9ab3531..a05c982b367 100644
--- a/intern/cycles/blender/blender_object.cpp
+++ b/intern/cycles/blender/blender_object.cpp
@@ -384,6 +384,23 @@ Object *BlenderSync::sync_object(BL::Object& b_parent,
 		object_updated = true;
 	}
 
+	/* sync the asset name for Cryptomatte */
+	BL::Object parent = b_ob.parent();
+	ustring parent_name;
+	if(parent) {
+		while(parent.parent()) {
+			parent = parent.parent();
+		}
+		parent_name = parent.name();
+	}
+	else {
+		parent_name = b_ob.name();
+	}
+	if(object->asset_name != parent_name) {
+		object->asset_name = parent_name;
+		object_updated = true;
+	}
+
 	/* object sync
 	 * transform comparison should not be needed, but duplis don't work perfect
 	 * in the depsgraph and may not signal changes, so this is a workaround */
diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp
index a07131d04ae..e05351eea40 100644
--- a/intern/cycles/blender/blender_session.cpp
+++ b/intern/cycles/blender/blender_session.cpp
@@ -405,17 +405,19 @@ void BlenderSession::render()
 		BL::RenderLayer b_rlay = *b_single_rlay;
 
 		/* add passes */
-		array<Pass> passes = sync->sync_render_passes(b_rlay, *b_layer_iter, session_params);
+		vector<Pass> passes = sync->sync_render_passes(b_rlay, *b_layer_iter, session_params);
 		buffer_params.passes = passes;
 
 		PointerRNA crl = RNA_pointer_get(&b_layer_iter->ptr, "cycles");
 		bool use_denoising = get_boolean(crl, "use_denoising");
+		bool denoising_passes = use_denoising || get_boolean(crl, "denoising_store_passes");
 
 		session->tile_manager.schedule_denoising = use_denoising;
-		buffer_params.denoising_data_pass = use_denoising;
+		buffer_params.denoising_data_pass = denoising_passes;
 		buffer_params.denoising_clean_pass = (scene->film->denoising_flags & DENOISING_CLEAN_ALL_PASSES);
 
 		session->params.use_denoising = use_denoising;
+		session->params.denoising_passes = denoising_passes;
 		session->params.denoising_radius = get_int(crl, "denoising_radius");
 		session->params.denoising_strength = get_float(crl, "denoising_strength");
 		session->params.denoising_feature_strength = get_float(crl, "denoising_feature_strength");
@@ -700,7 +702,7 @@ void BlenderSession::do_write_update_render_result(BL::RenderResult& b_rr,
 			bool read = false;
 			if(pass_type != PASS_NONE) {
 				/* copy pixels */
-				read = buffers->get_pass_rect(pass_type, exposure, sample, components, &pixels[0]);
+				read = buffers->get_pass_rect(pass_type, exposure, sample, components, &pixels[0], b_pass.name());
 			}
 			else {
 				int denoising_offset = BlenderSync::get_denoising_pass(b_pass);
@@ -719,7 +721,7 @@ void BlenderSession::do_write_update_render_result(BL::RenderResult& b_rr,
 	else {
 		/* copy combined pass */
 		BL::RenderPass b_combined_pass(b_rlay.passes.find_by_name("Combined", b_rview_name.c_str()));
-		if(buffers->get_pass_rect(PASS_COMBINED, exposure, sample, 4, &pixels[0]))
+		if(buffers->get_pass_rect(PASS_COMBINED, exposure, sample, 4, &pixels[0], "Combined"))
 			b_combined_pass.rect(&pixels[0]);
 	}
 
diff --git a/intern/cycles/blender/blender_sync.cpp b/intern/cycles/blender/blender_sync.cpp
index 3bb9b479fde..832847c179f 100644
--- a/intern/cycles/blender/blender_sync.cpp
+++ b/intern/cycles/blender/blender_sync.cpp
@@ -40,6 +40,8 @@
 
 CCL_NAMESPACE_BEGIN
 
+static const char *cryptomatte_prefix = "Crypto";
+
 /* Constructor */
 
 BlenderSync::BlenderSync(BL::RenderEngine& b_engine,
@@ -517,6 +519,9 @@ PassType BlenderSync::get_pass_type(BL::RenderPass& b_pass)
 	MAP_PASS("Debug Ray Bounces", PASS_RAY_BOUNCES);
 #endif
 	MAP_PASS("Debug Render Time", PASS_RENDER_TIME);
+	if(string_startswith(name, cryptomatte_prefix)) {
+		return PASS_CRYPTOMATTE;
+	}
 #undef MAP_PASS
 
 	return PASS_NONE;
@@ -525,6 +530,9 @@ PassType BlenderSync::get_pass_type(BL::RenderPass& b_pass)
 int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
 {
 	string name = b_pass.name();
+
+	if(name == "Noisy Image") return DENOISING_PASS_COLOR;
+
 	if(name.substr(0, 10) != "Denoising ") {
 		return -1;
 	}
@@ -539,7 +547,6 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
 	MAP_PASS("Depth Variance", DENOISING_PASS_DEPTH_VAR);
 	MAP_PASS("Shadow A", DENOISING_PASS_SHADOW_A);
 	MAP_PASS("Shadow B", DENOISING_PASS_SHADOW_B);
-	MAP_PASS("Image", DENOISING_PASS_COLOR);
 	MAP_PASS("Image Variance", DENOISING_PASS_COLOR_VAR);
 	MAP_PASS("Clean", DENOISING_PASS_CLEAN);
 #undef MAP_PASS
@@ -547,11 +554,11 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
 	return -1;
 }
 
-array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
-                                            BL::SceneRenderLayer& b_srlay,
-                                            const SessionParams &session_params)
+vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
+                                             BL::SceneRenderLayer& b_srlay,
+                                             const SessionParams &session_params)
 {
-	array<Pass> passes;
+	vector<Pass> passes;
 	Pass::add(PASS_COMBINED, passes);
 
 	if(!session_params.device.advanced_shading) {
@@ -571,22 +578,11 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
 			Pass::add(pass_type, passes);
 	}
 
-	scene->film->denoising_flags = 0;
 	PointerRNA crp = RNA_pointer_get(&b_srlay.ptr, "cycles");
-	if(get_boolean(crp, "denoising_store_passes") &&
-	   get_boolean(crp, "use_denoising"))
-	{
-		b_engine.add_pass("Denoising Normal",          3, "XYZ", b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Albedo",          3, "RGB", b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Albedo Variance", 3, "RGB", b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Depth",           1, "Z",   b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Depth Variance",  1, "Z",   b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Shadow A",        3, "XYV", b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Shadow B",        3, "XYV", b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Image",           3, "RGB", b_srlay.name().c_str());
-		b_engine.add_pass("Denoising Image Variance",  3, "RGB", b_srlay.name().c_str());
-
+	bool use_denoising = get_boolean(crp, "use_denoising");
+	bool store_denoising_passes = get_boolean(crp, "denoising_store_passes");
+	scene->film->denoising_flags = 0;
+	if(use_denoising || store_denoising_passes) {
 #define MAP_OPTION(name, flag) if(!get_boolean(crp, name)) scene->film->denoising_flags |= flag;
 		MAP_OPTION("denoising_diffuse_direct",        DENOISING_CLEAN_DIFFUSE_DIR);
 		MAP_OPTION("denoising_diffuse_indirect",      DENOISING_CLEAN_DIFFUSE_IND);
@@ -597,9 +593,22 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
 		MAP_OPTION("denoising_subsurface_direct",     DENOISING_CLEAN_SUBSURFACE_DIR);
 		MAP_OPTION("denoising_subsurface_indirect",   DENOISING_CLEAN_SUBSURFACE_IND);
 #undef MAP_OPTION
+		b_engine.add_pass("Noisy Image", 4, "RGBA", b_srlay.name().c_str());
+	}
+
+	if(store_denoising_passes) {
+		b_engine.add_pass("Denoising Normal",          3, "XYZ", b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Albedo",          3, "RGB", b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Albedo Variance", 3, "RGB", b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Depth",           1, "Z",   b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Depth Variance",  1, "Z",   b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Shadow A",        3, "XYV", b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Shadow B",        3, "XYV", b_srlay.name().c_str());
+		b_engine.add_pass("Denoising Image Variance",  3, "RGB", b_srlay.name().c_str());
 
 		if(scene->film->denoising_flags & DENOISING_CLEAN_ALL_PASSES) {
-			b_engine.add_pass("Denoising Clean", 3, "RGB", b_srlay.name().c_str());
+			b_engine.add_pass("Denoising Clean",   3, "RGB", b_srlay.name().c_str());
 		}
 	}
 #ifdef __KERNEL_DEBUG__
@@ -633,6 +642,39 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
 		Pass::add(PASS_VOLUME_INDIRECT, passes);
 	}
 
+	/* Cryptomatte stores two ID/weight pairs per RGBA layer.
+	 * User facing paramter is the number of pairs. */
+	int crypto_depth = min(16, get_int(crp, "pass_crypto_depth")) / 2;
+	scene->film->cryptomatte_depth = crypto_depth;
+	scene->film->cryptomatte_passes = CRYPT_NONE;
+	if(get_boolean(crp, "use_pass_crypto_object")) {
+		for(int i = 0; i < crypto_depth; ++i) {
+			string passname = cryptomatte_prefix + string_printf("Object%02d", i);
+			b_engine.add_pass(passname.c_str(), 4, "RGBA", b_srlay.name().c_str());
+			Pass::add(PASS_CRYPTOMATTE, passes, passname.c_str());
+		}
+		scene->film->cryptomatte_passes = (CryptomatteType)(scene->film->cryptomatte_passes | CRYPT_OBJECT);
+	}
+	if(get_boolean(crp, "use_pass_crypto_material")) {
+		for(int i = 0; i < crypto_depth; ++i) {
+			string passname = cryptomatte_prefix + string_printf("Material%02d", i);
+			b_engine.add_pass(passname.c_str(), 4, "RGBA", b_srlay.name().c_str());
+			Pass::add(PASS_CRYPTOMATTE, passes, passname.c_str());
+		}
+		scene->film->cryptomatte_passes = (CryptomatteType)(scene->film->cryptomatte_passes | CRYPT_MATERIAL);
+	}
+	if(get_boolean(crp, "use_pass_crypto_asset")) {
+		for(int i = 0; i < crypto_depth; ++i) {
+			string passname = cryptomatte_prefix + string_printf("Asset%02d", i);
+			b_engine.add_pass(passname.c_str(), 4, "RGBA", b_srlay.name().c_str());
+			Pass::add(PASS_CRYPTOMATTE, passes, passname.c_str());
+		}
+		scene->film->cryptomatte_passes = (CryptomatteType)(scene->film->cryptomatte_passes | CRYPT_ASSET);
+	}
+	if(get_boolean(crp, "pass_crypto_accurate") && scene->film->cryptomatte_passes != CRYPT_NONE) {
+		scene->film->cryptomatte_passes = (CryptomatteType)(scene->film->cryptomatte_passes | CRYPT_ACCURATE);
+	}
+
 	return passes;
 }
 
diff --git a/intern/cycles/blender/blender_sync.h b/intern/cycles/blender/blender_sync.h
index 5e63f76033d..eb84bedb118 100644
--- a/intern/cycles/blender/blender_sync.h
+++ b/intern/cycles/blender/blender_sync.h
@@ -66,9 +66,9 @@ public:
 	               void **python_thread_state,
 	               const char *layer = 0);
 	void sync_render_layers(BL::SpaceView3D& b_v3d, const char *layer);
-	array<Pass> sync_render_passes(BL::RenderLayer& b_rlay,
-	                               BL::SceneRenderLayer& b_srlay,
-	                               const SessionParams &session_params);
+	vector<Pass> sync_render_passes(BL::RenderLayer& b_rlay,
+	                                BL::SceneRenderLayer& b_srlay,
+	                                const SessionParams &session_params);
 	void sync_integrator();
 	void sync_camera(BL::RenderSettings& b_render,
 	                 BL::Object& b_override,
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp
index 7322086e655..4db626b6bfd 100644
--- a/intern/cycles/device/device_cpu.cpp
+++ b/intern/cycles/device/device_cpu.cpp
@@ -41,6 +41,7 @@
 #include "kernel/osl/osl_globals.h"
 
 #include "render/buffers.h"
+#include "render/coverage.h"
 
 #include "util/util_debug.h"
 #include "util/util_foreach.h"
@@ -677,8 +678,15 @@ public:
 
 	void path_trace(DeviceTask &task, RenderTile &tile, KernelGlobals *kg)
 	{
+		const bool use_coverage = kernel_data.film.cryptomatte_passes & CRYPT_ACCURATE;
+
 		scoped_timer timer(&tile.buffers->render_time);
 
+		Coverage coverage(kg, tile);
+		if(use_coverage) {
+			coverage.init_path_trace();
+		}
+
 		float *render_buffer = (float*)tile.buffer;
 		int start_sample = tile.start_sample;
 		int end_sample = tile.start_sample + tile.num_samples;
@@ -694,6 +702,9 @@ public:
 
 			for(int y = tile.y; y < tile.y + tile.h; y++) {
 				for(int x = tile.x; x < tile.x + tile.w; x++) {
+					if(use_coverage) {
+						coverage.init_pixel(x, y);
+					}
 					path_trace_kernel()(kg, render_buffer,
 					                    sample, x, y, tile.offset, tile.stride);
 				}
@@ -703,6 +714,9 @@ public:
 
 			task.update_progress(&tile, tile.w*tile.h);
 		}
+		if(use_coverage) {
+			coverage.finalize();
+		}
 	}
 
 	void denoise(DenoisingTask& denoising, RenderTile &tile)
@@ -763,7 +777,6 @@ public:
 			}
 			else if(tile.task == RenderTile::DENOISE) {
 				denoise(denoising, tile);
-
 				task.update_progress(&tile, tile.w*tile.h);
 			}
 
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 64b43a3c626..92cb66bdec9 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -97,6 +97,7 @@ set(SRC_HEADERS
 	kernel_emission.h
 	kernel_film.h
 	kernel_globals.h
+	kernel_id_passes.h
 	kernel_jitter.h
 	kernel_light.h
 	kernel_math.h
diff --git a/intern/cycles/kernel/geom/geom_object.h b/intern/cycles/kernel/geom/geom_object.h
index 60bc30d2e2a..669c932d720 100644
--- a/intern/cycles/kernel/geom/geom_object.h
+++ b/intern/cycles/kernel/geom/geom_object.h
@@ -310,6 +310,24 @@ ccl_device int shader_pass_id(KernelGlobals *kg, const ShaderData *sd)
 	return kernel_tex_fetch(__shaders, (sd->shader & SHADER_MASK)).pass_id;
 }
 
+/* Cryptomatte ID */
+
+ccl_device_inline float object_cryptomatte_id(KernelGlobals *kg, int object)
+{
+	if(object == OBJECT_NONE)
+		return 0.0f;
+
+	return kernel_tex_fetch(__objects, object).cryptomatte_object;
+}
+
+ccl_device_inline float object_cryptomatte_asset_id(KernelGlobals *kg, int object)
+{
+	if(object == OBJECT_NONE)
+		return 0;
+
+	return kernel_tex_fetch(__objects, object).cryptomatte_asset;
+}
+
 /* Particle data from which object was instanced */
 
 ccl_device_inline uint particle_index(KernelGlobals *kg, int particle)
diff --git a/intern/cycles/kernel/kernel_globals.h b/intern/cycles/kernel/kernel_globals.h
index 74cfacb5bc1..37402f42863 100644
--- a/intern/cycles/kernel/kernel_globals.h
+++ b/intern/cycles/kernel/kernel_globals.h
@@ -21,6 +21,7 @@
 
 #ifdef __KERNEL_CPU__
 #  include "util/util_vector.h"
+#  include "util/util_map.h"
 #endif
 
 #ifdef __KERNEL_OPENCL__
@@ -42,6 +43,8 @@ struct OSLThreadData;
 struct OSLShadingSystem;
 #  endif
 
+typedef unordered_map<float, float> CoverageMap;
+
 struct Intersection;
 struct VolumeStep;
 
@@ -68,6 +71,11 @@ typedef struct KernelGlobals {
 	VolumeStep *decoupled_volume_steps[2];
 	int decoupled_volume_steps_index;
 
+	/* A buffer for storing per-pixel coverage for Cryptomatte. */
+	CoverageMap *coverage_object;
+	CoverageMap *coverage_material;
+	CoverageMap *coverage_asset;
+
 	/* split kernel */
 	SplitData split_data;
 	SplitParams split_param_data;
diff --git a/intern/cycles/kernel/kernel_id_passes.h b/intern/cycles/kernel/kernel_id_passes.h
new file mode 100644
index 00000000000..486c61d2ae5
--- /dev/null
+++ b/intern/cycles/kernel/kernel_id_passes.h
@@ -0,0 +1,94 @@
+/*
+* Copyright 2018 Blender Foundation
+*
+* Licensed under the Apache License, Version 2.0 (the "License");
+* you may not use this file except in compliance with the License.
+* You may obtain a copy of the License at
+*
+* http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing, software
+* distributed under the License is distributed on an "AS IS" BASIS,
+* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+* See the License for the specific language governing permissions and
+* limitations under the License.
+*/
+
+CCL_NAMESPACE_BEGIN
+
+ccl_device_inline void kernel_write_id_slots(ccl_global float *buffer, int num_slots, float id, float weight)
+{
+	kernel_assert(id != ID_NONE);
+	if(weight == 0.0f) {
+		return;
+	}
+	
+	for(int slot = 0; slot < num_slots; slot++) {
+		ccl_global float2 *id_buffer = (ccl_global float2*)buffer;
+#ifdef __ATOMIC_PASS_WRITE__
+		/* If the loop reaches an empty slot, the ID isn't in any slot yet - so add it! */
+		if(id_buffer[slot].x == ID_NONE) {
+			/* Use an atomic to claim this slot.
+			* If a different thread got here first, try again from this slot on. */
+			float old_id = atomic_compare_and_swap_float(buffer+slot*2, ID_NONE, id);
+			if(old_id != ID_NONE && old_id != id) {
+				continue;
+			}
+			atomic_add_and_fetch_float(buffer+slot*2+1, weight);
+			break;
+		}
+		/* If there already is a slot for that ID, add the weight.
+		 * If no slot was found, add it to the last. */
+		else if(id_buffer[slot].x == id || slot == num_slots - 1) {
+			atomic_add_and_fetch_float(buffer+slot*2+1, weight);
+			break;
+		}
+#else /* __ATOMIC_PASS_WRITE__ */
+		/* If the loop reaches an empty slot, the ID isn't in any slot yet - so add it! */
+		if(id_buffer[slot].x == ID_NONE) {
+			id_buffer[slot].x = id;
+			id_buffer[slot].y = weight;
+			break;
+		}
+		/* If there already is a slot for that ID, add the weight.
+		* If no slot was found, add it to the last. */
+		else if(id_buffer[slot].x == id || slot == num_slots - 1) {
+			id_buffer[slot].y += weight;
+			break;
+		}
+#endif /* __ATOMIC_PASS_WRITE__ */
+	}
+}
+
+ccl_device_inline void kernel_sort_id_slots(ccl_global float *buffer, int num_slots)
+{
+	ccl_global float2 *id_buffer = (ccl_global float2*)buffer;
+	for(int slot = 1; slot < num_slots; ++slot) {
+		if(id_buffer[slot].x == ID_NONE) {
+			return;
+		}
+		/* Since we're dealing with a tiny number of elements, insertion sort should be fine. */
+		int i = slot;
+		while(i > 0 && id_buffer[i].y > id_buffer[i - 1].y) {
+			float2 swap = id_buffer[i];
+			id_buffer[i] = id_buffer[i - 1];
+			id_buffer[i - 1] = swap;
+			--i;
+		}
+	}
+}
+
+#ifdef __KERNEL_GPU__
+/* post-sorting for Cryptomatte */
+ccl_device void kernel_cryptomatte_post(KernelGlobals *kg, ccl_global float *buffer, uint sample, int x, int y, int offset, int stride)
+{
+	if(sample - 1 == kernel_data.integrator.aa_samples) {
+		int index = offset + x + y * stride;
+		int pass_stride = kernel_data.film.pass_stride;
+		ccl_global float *cryptomatte_buffer = buffer + index * pass_stride + kernel_data.film.pass_cryptomatte;
+		kernel_sort_id_slots(cryptomatte_buffer, 2 * kernel_data.film.cryptomatte_depth);
+	}
+}
+#endif
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_montecarlo.h b/intern/cycles/kernel/kernel_montecarlo.h
index 9b96bb80c32..ce68aa16af8 100644
--- a/intern/cycles/kernel/kernel_montecarlo.h
+++ b/intern/cycles/kernel/kernel_montecarlo.h
@@ -187,7 +187,10 @@ ccl_device float2 regular_polygon_sample(float corners, float rotation, float u,
 ccl_device float3 ensure_valid_reflection(float3 Ng, float3 I, float3 N)
 {
 	float3 R = 2*dot(N, I)*N - I;
-	if(dot(Ng, R) >= 0.05f) {
+
+	/* Reflection rays may always be at least as shallow as the incoming ray. */
+	float threshold = min(0.9f*dot(Ng, I), 0.01f);
+	if(dot(Ng, R) >= threshold) {
 		return N;
 	}
 
@@ -195,22 +198,86 @@ ccl_device float3 ensure_valid_reflection(float3 Ng, float3 I, float3 N)
 	 * The X axis is found by normalizing the component of N that's orthogonal to Ng.
 	 * The Y axis isn't actually needed.
 	 */
-	float3 X = normalize(N - dot(N, Ng)*Ng);
-
-	/* Calculate N.z and N.x in the local coordinate system. */
-	float Iz = dot(I, Ng);
-	float Ix2 = sqr(dot(I, X)), Iz2 = sqr(Iz);
-	float Ix2Iz2 = Ix2 + Iz2;
-
-	float a = safe_sqrtf(Ix2*(Ix2Iz2 - sqr(0.05f)));
-	float b = Iz*0.05f + Ix2Iz2;
-	float c = (a + b > 0.0f)? (a + b) : (-a + b);
+	float NdotNg = dot(N, Ng);
+	float3 X = normalize(N - NdotNg*Ng);
+
+	/* Calculate N.z and N.x in the local coordinate system.
+	 *
+	 * The goal of this computation is to find a N' that is rotated towards Ng just enough
+	 * to lift R' above the threshold (here called t), therefore dot(R', Ng) = t.
+	 *
+	 * According to the standard reflection equation, this means that we want dot(2*dot(N', I)*N' - I, Ng) = t.
+	 *
+	 * Since the Z axis of our local coordinate system is Ng, dot(x, Ng) is just x.z, so we get 2*dot(N', I)*N'.z - I.z = t.
+	 *
+	 * The rotation is simple to express in the coordinate system we formed - since N lies in the X-Z-plane, we know that
+	 * N' will also lie in the X-Z-plane, so N'.y = 0 and therefore dot(N', I) = N'.x*I.x + N'.z*I.z .
+	 *
+	 * Furthermore, we want N' to be normalized, so N'.x = sqrt(1 - N'.z^2).
+	 *
+	 * With these simplifications, we get the final equation 2*(sqrt(1 - N'.z^2)*I.x + N'.z*I.z)*N'.z - I.z = t.
+	 *
+	 * The only unknown here is N'.z, so we can solve for that.
+	 *
+	 * The equation has four solutions in general:
+	 *
+	 * N'.z = +-sqrt(0.5*(+-sqrt(I.x^2*(I.x^2 + I.z^2 - t^2)) + t*I.z + I.x^2 + I.z^2)/(I.x^2 + I.z^2))
+	 * We can simplify this expression a bit by grouping terms:
+	 *
+	 * a = I.x^2 + I.z^2
+	 * b = sqrt(I.x^2 * (a - t^2))
+	 * c = I.z*t + a
+	 * N'.z = +-sqrt(0.5*(+-b + c)/a)
+	 *
+	 * Two solutions can immediately be discarded because they're negative so N' would lie in the lower hemisphere.
+	 */
+	float Ix = dot(I, X), Iz = dot(I, Ng);
+	float Ix2 = sqr(Ix), Iz2 = sqr(Iz);
+	float a = Ix2 + Iz2;
+
+	float b = safe_sqrtf(Ix2*(a - sqr(threshold)));
+	float c = Iz*threshold + a;
+
+	/* Evaluate both solutions.
+	 * In many cases one can be immediately discarded (if N'.z would be imaginary or larger than one), so check for that first.
+	 * If no option is viable (might happen in extreme cases like N being in the wrong hemisphere), give up and return Ng. */
+	float fac = 0.5f/a;
+	float N1_z2 = fac*(b+c), N2_z2 = fac*(-b+c);
+	bool valid1 = (N1_z2 > 1e-5f) && (N1_z2 <= (1.0f + 1e-5f));
+	bool valid2 = (N2_z2 > 1e-5f) && (N2_z2 <= (1.0f + 1e-5f));
+
+	float2 N_new;
+	if(valid1 && valid2) {
+		/* If both are possible, do the expensive reflection-based check. */
+		float2 N1 = make_float2(safe_sqrtf(1.0f - N1_z2), safe_sqrtf(N1_z2));
+		float2 N2 = make_float2(safe_sqrtf(1.0f - N2_z2), safe_sqrtf(N2_z2));
+
+		float R1 = 2*(N1.x*Ix + N1.y*Iz)*N1.y - Iz;
+		float R2 = 2*(N2.x*Ix + N2.y*Iz)*N2.y - Iz;
+
+		valid1 = (R1 >= 1e-5f);
+		valid2 = (R2 >= 1e-5f);
+		if(valid1 && valid2) {
+			/* If both solutions are valid, return the one with the shallower reflection since it will be closer to the input
+			 * (if the original reflection wasn't shallow, we would not be in this part of the function). */
+			N_new = (R1 < R2)? N1 : N2;
+		}
+		else {
+			/* If only one reflection is valid (= positive), pick that one. */
+			N_new = (R1 > R2)? N1 : N2;
+		}
 
-	float Nz = safe_sqrtf(0.5f * c * (1.0f / Ix2Iz2));
-	float Nx = safe_sqrtf(1.0f - sqr(Nz));
+	}
+	else if(valid1 || valid2) {
+		/* Only one solution passes the N'.z criterium, so pick that one. */
+		float Nz2 = valid1? N1_z2 : N2_z2;
+		N_new = make_float2(safe_sqrtf(1.0f - Nz2), safe_sqrtf(Nz2));
+	}
+	else {
+		return Ng;
+	}
 
-	/* Transform back into global coordinates. */
-	return Nx*X + Nz*Ng;
+	return N_new.x*X + N_new.y*Ng;
 }
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_passes.h b/intern/cycles/kernel/kernel_passes.h
index 458aa6c2a97..e256a1819ed 100644
--- a/intern/cycles/kernel/kernel_passes.h
+++ b/intern/cycles/kernel/kernel_passes.h
@@ -14,12 +14,14 @@
  * limitations under the License.
  */
 
-CCL_NAMESPACE_BEGIN
-
 #if defined(__SPLIT_KERNEL__) || defined(__KERNEL_CUDA__)
 #define __ATOMIC_PASS_WRITE__
 #endif
 
+#include "kernel/kernel_id_passes.h"
+
+CCL_NAMESPACE_BEGIN
+
 ccl_device_inline void kernel_write_pass_float(ccl_global float *buffer, float value)
 {
 	ccl_global float *buf = buffer;
@@ -189,6 +191,23 @@ ccl_device_inline void kernel_write_debug_passes(KernelGlobals *kg,
 }
 #endif /* __KERNEL_DEBUG__ */
 
+#ifdef __KERNEL_CPU__
+#define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name) kernel_write_id_pass_cpu(buffer, depth * 2, id, matte_weight, kg->coverage_##name)
+ccl_device_inline size_t kernel_write_id_pass_cpu(float *buffer, size_t depth, float id, float matte_weight, CoverageMap *map)
+{
+	if(map) {
+		(*map)[id] += matte_weight;
+		return 0;
+	}
+#else /* __KERNEL_CPU__ */
+#define WRITE_ID_SLOT(buffer, depth, id, matte_weight, name) kernel_write_id_slots_gpu(buffer, depth * 2, id, matte_weight) 
+ccl_device_inline size_t kernel_write_id_slots_gpu(ccl_global float *buffer, size_t depth, float id, float matte_weight)
+{
+#endif /* __KERNEL_CPU__ */
+	kernel_write_id_slots(buffer, depth, id, matte_weight);
+	return depth * 2;
+}
+
 ccl_device_inline void kernel_write_data_passes(KernelGlobals *kg, ccl_global float *buffer, PathRadiance *L,
 	ShaderData *sd, ccl_addr_space PathState *state, float3 throughput)
 {
@@ -242,6 +261,26 @@ ccl_device_inline void kernel_write_data_passes(KernelGlobals *kg, ccl_global fl
 		}
 	}
 
+	if(kernel_data.film.cryptomatte_passes) {
+		const float matte_weight = average(throughput) * (1.0f - average(shader_bsdf_transparency(kg, sd)));
+		if(matte_weight > 0.0f) {
+			ccl_global float *cryptomatte_buffer = buffer + kernel_data.film.pass_cryptomatte;
+			if(kernel_data.film.cryptomatte_passes & CRYPT_OBJECT) {
+				float id = object_cryptomatte_id(kg, sd->object);
+				cryptomatte_buffer += WRITE_ID_SLOT(cryptomatte_buffer, kernel_data.film.cryptomatte_depth, id, matte_weight, object);
+			}
+			if(kernel_data.film.cryptomatte_passes & CRYPT_MATERIAL) {
+				float id = shader_cryptomatte_id(kg, sd->shader);
+				cryptomatte_buffer += WRITE_ID_SLOT(cryptomatte_buffer, kernel_data.film.cryptomatte_depth, id, matte_weight, material);
+			}
+			if(kernel_data.film.cryptomatte_passes & CRYPT_ASSET) {
+				float id = object_cryptomatte_asset_id(kg, sd->object);
+				cryptomatte_buffer += WRITE_ID_SLOT(cryptomatte_buffer, kernel_data.film.cryptomatte_depth, id, matte_weight, asset);
+			}
+		}
+	}
+
+
 	if(light_flag & PASSMASK_COMPONENT(DIFFUSE))
 		L->color_diffuse += shader_bsdf_diffuse(kg, sd)*throughput;
 	if(light_flag & PASSMASK_COMPONENT(GLOSSY))
diff --git a/intern/cycles/kernel/kernel_shader.h b/intern/cycles/kernel/kernel_shader.h
index e834b701f96..af883aa715b 100644
--- a/intern/cycles/kernel/kernel_shader.h
+++ b/intern/cycles/kernel/kernel_shader.h
@@ -1276,4 +1276,9 @@ ccl_device bool shader_transparent_shadow(KernelGlobals *kg, Intersection *isect
 }
 #endif  /* __TRANSPARENT_SHADOWS__ */
 
+ccl_device float shader_cryptomatte_id(KernelGlobals *kg, int shader)
+{
+	return kernel_tex_fetch(__shaders, (shader & SHADER_MASK)).cryptomatte_id;
+}
+
 CCL_NAMESPACE_END
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index f799dff3a43..04b07d4dd5e 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -58,6 +58,7 @@ CCL_NAMESPACE_BEGIN
 #define OBJECT_NONE				(~0)
 #define PRIM_NONE				(~0)
 #define LAMP_NONE				(~0)
+#define ID_NONE					(0.0f)
 
 #define VOLUME_STACK_SIZE		32
 
@@ -423,6 +424,7 @@ typedef enum PassType {
 	PASS_RAY_BOUNCES,
 #endif
 	PASS_RENDER_TIME,
+	PASS_CRYPTOMATTE,
 	PASS_CATEGORY_MAIN_END = 31,
 
 	PASS_MIST = 32,
@@ -451,6 +453,14 @@ typedef enum PassType {
 
 #define PASS_ANY (~0)
 
+typedef enum CryptomatteType {
+	CRYPT_NONE = 0,
+	CRYPT_OBJECT = (1 << 0),
+	CRYPT_MATERIAL = (1 << 1),
+	CRYPT_ASSET = (1 << 2),
+	CRYPT_ACCURATE = (1 << 3),
+} CryptomatteType;
+
 typedef enum DenoisingPassOffsets {
 	DENOISING_PASS_NORMAL             = 0,
 	DENOISING_PASS_NORMAL_VAR         = 3,
@@ -1271,17 +1281,20 @@ typedef struct KernelFilm {
 	int pass_shadow;
 	float pass_shadow_scale;
 	int filter_table_offset;
+	int cryptomatte_passes;
+	int cryptomatte_depth;
+	int pass_cryptomatte;
 
 	int pass_mist;
 	float mist_start;
 	float mist_inv_depth;
 	float mist_falloff;
-
+	
 	int pass_denoising_data;
 	int pass_denoising_clean;
 	int denoising_flags;
 
-	int pad1, pad2, pad3;
+	int pad1, pad2;
 
 	/* XYZ to rendering color space transform. float4 instead of float3 to
 	 * ensure consistent padding/alignment across devices. */
@@ -1477,7 +1490,11 @@ typedef struct KernelObject {
 	uint patch_map_offset;
 	uint attribute_map_offset;
 	uint motion_offset;
-	uint pad;
+	uint pad1;
+
+	float cryptomatte_object;
+	float cryptomatte_asset;
+	float pad2, pad3;
 } KernelObject;
 static_assert_align(KernelObject, 16);
 
@@ -1557,7 +1574,7 @@ static_assert_align(KernelParticle, 16);
 
 typedef struct KernelShader {
 	float constant_emission[3];
-	float pad1;
+	float cryptomatte_id;
 	int flags;
 	int pass_id;
 	int pad2, pad3;
diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu
index 8a180a509e8..af311027f78 100644
--- a/intern/cycles/kernel/kernels/cuda/kernel.cu
+++ b/intern/cycles/kernel/kernels/cuda/kernel.cu
@@ -40,14 +40,21 @@ CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
 kernel_cuda_path_trace(WorkTile *tile, uint total_work_size)
 {
 	int work_index = ccl_global_id(0);
-
-	if(work_index < total_work_size) {
-		uint x, y, sample;
+	bool thread_is_active = work_index < total_work_size;
+	uint x, y, sample;
+	KernelGlobals kg;
+	if(thread_is_active) {
 		get_work_pixel(tile, work_index, &x, &y, &sample);
 
-		KernelGlobals kg;
 		kernel_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
 	}
+
+	if(kernel_data.film.cryptomatte_passes) {
+		__syncthreads();
+		if(thread_is_active) {
+			kernel_cryptomatte_post(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
+		}
+	}
 }
 
 #ifdef __BRANCHED_PATH__
@@ -56,14 +63,21 @@ CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_BRANCHED_MAX_REGISTERS)
 kernel_cuda_branched_path_trace(WorkTile *tile, uint total_work_size)
 {
 	int work_index = ccl_global_id(0);
-
-	if(work_index < total_work_size) {
-		uint x, y, sample;
+	bool thread_is_active = work_index < total_work_size;
+	uint x, y, sample;
+	KernelGlobals kg;
+	if(thread_is_active) {
 		get_work_pixel(tile, work_index, &x, &y, &sample);
 
-		KernelGlobals kg;
 		kernel_branched_path_trace(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
 	}
+	
+	if(kernel_data.film.cryptomatte_passes) {
+		__syncthreads();
+		if(thread_is_active) {
+			kernel_cryptomatte_post(&kg, tile->buffer, sample, x, y, tile->offset, tile->stride);
+		}
+	}
 }
 #endif
 
diff --git a/intern/cycles/kernel/kernels/opencl/kernel.cl b/intern/cycles/kernel/kernels/opencl/kernel.cl
index 63128d0aecf..de1f5088629 100644
--- a/intern/cycles/kernel/kernels/opencl/kernel.cl
+++ b/intern/cycles/kernel/kernels/opencl/kernel.cl
@@ -66,9 +66,17 @@ __kernel void kernel_ocl_path_trace(
 
 	int x = sx + ccl_global_id(0);
 	int y = sy + ccl_global_id(1);
-
-	if(x < sx + sw && y < sy + sh)
+	bool thread_is_active = x < sx + sw && y < sy + sh;
+	if(thread_is_active) {
 		kernel_path_trace(kg, buffer, sample, x, y, offset, stride);
+	}
+	if(kernel_data.film.cryptomatte_passes) {
+		/* Make sure no thread is writing to the buffers. */
+		ccl_barrier(CCL_LOCAL_MEM_FENCE);
+		if(thread_is_active) {
+			kernel_cryptomatte_post(kg, buffer, sample, x, y, offset, stride);
+		}
+	}
 }
 
 #else  /* __COMPILE_ONLY_MEGAKERNEL__ */
diff --git a/intern/cycles/kernel/osl/osl_services.cpp b/intern/cycles/kernel/osl/osl_services.cpp
index 7902381440b..81348f5594d 100644
--- a/intern/cycles/kernel/osl/osl_services.cpp
+++ b/intern/cycles/kernel/osl/osl_services.cpp
@@ -884,6 +884,23 @@ bool OSLRenderServices::has_userdata(ustring name, TypeDesc type, OSL::ShaderGlo
 	return false; /* never called by OSL */
 }
 
+TextureSystem::TextureHandle *OSLRenderServices::get_texture_handle(ustring filename)
+{
+	if (filename.length() && filename[0] == '@') {
+		/* Dummy, we don't use texture handles for builtin textures but need
+		 * to tell the OSL runtime optimizer that this is a valid texture. */
+		return NULL;
+	}
+	else {
+		return texturesys()->get_texture_handle(filename);
+	}
+}
+
+bool OSLRenderServices::good(TextureSystem::TextureHandle *texture_handle)
+{
+	return texturesys()->good(texture_handle);
+}
+
 bool OSLRenderServices::texture(ustring filename,
                                 TextureHandle *texture_handle,
                                 TexturePerthread *texture_thread_info,
@@ -894,7 +911,8 @@ bool OSLRenderServices::texture(ustring filename,
                                 int nchannels,
                                 float *result,
                                 float *dresultds,
-                                float *dresultdt)
+                                float *dresultdt,
+                                ustring *errormessage)
 {
 	OSL::TextureSystem *ts = osl_ts;
 	ShaderData *sd = (ShaderData *)(sg->renderstate);
@@ -1156,7 +1174,13 @@ bool OSLRenderServices::get_texture_info(OSL::ShaderGlobals *sg, ustring filenam
                                          TypeDesc datatype, void *data)
 {
 	OSL::TextureSystem *ts = osl_ts;
-	return ts->get_texture_info(filename, subimage, dataname, datatype, data);
+	if (filename.length() && filename[0] == '@') {
+		/* Special builtin textures. */
+		return false;
+	}
+	else {
+		return ts->get_texture_info(filename, subimage, dataname, datatype, data);
+	}
 }
 
 int OSLRenderServices::pointcloud_search(OSL::ShaderGlobals *sg, ustring filename, const OSL::Vec3 &center,
diff --git a/intern/cycles/kernel/osl/osl_services.h b/intern/cycles/kernel/osl/osl_services.h
index 50044746fd1..5dcaa4d7445 100644
--- a/intern/cycles/kernel/osl/osl_services.h
+++ b/intern/cycles/kernel/osl/osl_services.h
@@ -93,6 +93,10 @@ public:
 	bool getmessage(OSL::ShaderGlobals *sg, ustring source, ustring name,
 	                TypeDesc type, void *val, bool derivatives);
 
+	TextureSystem::TextureHandle *get_texture_handle(ustring filename);
+
+	bool good(TextureSystem::TextureHandle *texture_handle);
+
 	bool texture(ustring filename,
 	             TextureSystem::TextureHandle *texture_handle,
 	             TexturePerthread *texture_thread_info,
@@ -103,7 +107,8 @@ public:
 	             int nchannels,
 	             float *result,
 	             float *dresultds,
-	             float *dresultdt);
+	             float *dresultdt,
+	             ustring *errormessage);
 
 	bool texture3d(ustring filename,
 	               TextureHandle *texture_handle,
diff --git a/intern/cycles/kernel/shaders/stdosl.h b/intern/cycles/kernel/shaders/stdosl.h
index 4a8378796ba..f1235500f2b 100644
--- a/intern/cycles/kernel/shaders/stdosl.h
+++ b/intern/cycles/kernel/shaders/stdosl.h
@@ -284,33 +284,63 @@ point rotate (point p, float angle, point a, point b)
 
 normal ensure_valid_reflection(normal Ng, vector I, normal N)
 {
+    /* The implementation here mirrors the one in kernel_montecarlo.h,
+     * check there for an explanation of the algorithm. */
+
     float sqr(float x) { return x*x; }
 
     vector R = 2*dot(N, I)*N - I;
-    if (dot(Ng, R) >= 0.05) {
+
+    float threshold = min(0.9*dot(Ng, I), 0.01);
+    if(dot(Ng, R) >= threshold) {
         return N;
     }
 
-    /* Form coordinate system with Ng as the Z axis and N inside the X-Z-plane.
-     * The X axis is found by normalizing the component of N that's orthogonal to Ng.
-     * The Y axis isn't actually needed.
-     */
-    vector X = normalize(N - dot(N, Ng)*Ng);
+    float NdotNg = dot(N, Ng);
+    vector X = normalize(N - NdotNg*Ng);
 
-    /* Calculate N.z and N.x in the local coordinate system. */
     float Ix = dot(I, X), Iz = dot(I, Ng);
-    float Ix2 = sqr(dot(I, X)), Iz2 = sqr(dot(I, Ng));
-    float Ix2Iz2 = Ix2 + Iz2;
-
-    float a = sqrt(Ix2*(Ix2Iz2 - sqr(0.05)));
-    float b = Iz*0.05 + Ix2Iz2;
-    float c = (a + b > 0.0)? (a + b) : (-a + b);
+    float Ix2 = sqr(Ix), Iz2 = sqr(Iz);
+    float a = Ix2 + Iz2;
+
+    float b = sqrt(Ix2*(a - sqr(threshold)));
+    float c = Iz*threshold + a;
+
+    float fac = 0.5/a;
+    float N1_z2 = fac*(b+c), N2_z2 = fac*(-b+c);
+    int valid1 = (N1_z2 > 1e-5) && (N1_z2 <= (1.0 + 1e-5));
+    int valid2 = (N2_z2 > 1e-5) && (N2_z2 <= (1.0 + 1e-5));
+
+    float N_new_x, N_new_z;
+    if(valid1 && valid2) {
+        float N1_x = sqrt(1.0 - N1_z2), N1_z = sqrt(N1_z2);
+        float N2_x = sqrt(1.0 - N2_z2), N2_z = sqrt(N2_z2);
+
+        float R1 = 2*(N1_x*Ix + N1_z*Iz)*N1_z - Iz;
+        float R2 = 2*(N2_x*Ix + N2_z*Iz)*N2_z - Iz;
+
+        valid1 = (R1 >= 1e-5);
+        valid2 = (R2 >= 1e-5);
+        if(valid1 && valid2) {
+            N_new_x = (R1 < R2)? N1_x : N2_x;
+            N_new_z = (R1 < R2)? N1_z : N2_z;
+        }
+        else {
+            N_new_x = (R1 > R2)? N1_x : N2_x;
+            N_new_z = (R1 > R2)? N1_z : N2_z;
+        }
 
-    float Nz = sqrt(0.5 * c * (1.0 / Ix2Iz2));
-    float Nx = sqrt(1.0 - sqr(Nz));
+    }
+    else if(valid1 || valid2) {
+        float Nz2 = valid1? N1_z2 : N2_z2;
+        N_new_x = sqrt(1.0 - Nz2);
+        N_new_z = sqrt(Nz2);
+    }
+    else {
+        return Ng;
+    }
 
-    /* Transform back into global coordinates. */
-    return Nx*X + Nz*Ng;
+    return N_new_x*X + N_new_z*Ng;
 }
 
 
diff --git a/intern/cycles/kernel/split/kernel_buffer_update.h b/intern/cycles/kernel/split/kernel_buffer_update.h
index 180c0b57077..18eec6372f1 100644
--- a/intern/cycles/kernel/split/kernel_buffer_update.h
+++ b/intern/cycles/kernel/split/kernel_buffer_update.h
@@ -80,8 +80,10 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg,
 	PathRadiance *L = &kernel_split_state.path_radiance[ray_index];
 	ccl_global Ray *ray = &kernel_split_state.ray[ray_index];
 	ccl_global float3 *throughput = &kernel_split_state.throughput[ray_index];
+	bool ray_was_updated = false;
 
 	if(IS_STATE(ray_state, ray_index, RAY_UPDATE_BUFFER)) {
+		ray_was_updated = true;
 		uint sample = state->sample;
 		uint buffer_offset = kernel_split_state.buffer_offset[ray_index];
 		ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset;
@@ -92,6 +94,17 @@ ccl_device void kernel_buffer_update(KernelGlobals *kg,
 		ASSIGN_RAY_STATE(ray_state, ray_index, RAY_TO_REGENERATE);
 	}
 
+	if(kernel_data.film.cryptomatte_passes) {
+		/* Make sure no thread is writing to the buffers. */
+		ccl_barrier(CCL_LOCAL_MEM_FENCE);
+		if(ray_was_updated && state->sample - 1 == kernel_data.integrator.aa_samples) {
+			uint buffer_offset = kernel_split_state.buffer_offset[ray_index];
+			ccl_global float *buffer = kernel_split_params.tile.buffer + buffer_offset;
+			ccl_global float *cryptomatte_buffer = buffer + kernel_data.film.pass_cryptomatte;
+			kernel_sort_id_slots(cryptomatte_buffer, 2 * kernel_data.film.cryptomatte_depth);
+		}
+	}
+
 	if(IS_STATE(ray_state, ray_index, RAY_TO_REGENERATE)) {
 		/* We have completed current work; So get next work */
 		ccl_global uint *work_pools = kernel_split_params.work_pools;
diff --git a/intern/cycles/render/CMakeLists.txt b/intern/cycles/render/CMakeLists.txt
index 7d2220f37f9..c0ce7368771 100644
--- a/intern/cycles/render/CMakeLists.txt
+++ b/intern/cycles/render/CMakeLists.txt
@@ -15,6 +15,7 @@ set(SRC
 	buffers.cpp
 	camera.cpp
 	constant_fold.cpp
+	coverage.cpp
 	film.cpp
 	graph.cpp
 	image.cpp
@@ -46,6 +47,7 @@ set(SRC_HEADERS
 	buffers.h
 	camera.h
 	constant_fold.h
+	coverage.h
 	film.h
 	graph.h
 	image.h
diff --git a/intern/cycles/render/buffers.cpp b/intern/cycles/render/buffers.cpp
index 4cd8b3726d3..dd20efb3dde 100644
--- a/intern/cycles/render/buffers.cpp
+++ b/intern/cycles/render/buffers.cpp
@@ -160,11 +160,12 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
 	                (offset == DENOISING_PASS_DEPTH_VAR) ||
 	                (offset == DENOISING_PASS_COLOR_VAR);
 
+	float scale_exposure = scale;
 	if(offset == DENOISING_PASS_COLOR || offset == DENOISING_PASS_CLEAN) {
-		scale *= exposure;
+		scale_exposure *= exposure;
 	}
 	else if(offset == DENOISING_PASS_COLOR_VAR) {
-		scale *= exposure*exposure;
+		scale_exposure *= exposure*exposure;
 	}
 
 	offset += params.get_denoising_offset();
@@ -181,14 +182,14 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
 
 		if(components == 1) {
 			for(int i = 0; i < size; i++, mean += pass_stride, var += pass_stride, pixels++) {
-				pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale;
+				pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale_exposure;
 			}
 		}
 		else if(components == 3) {
 			for(int i = 0; i < size; i++, mean += pass_stride, var += pass_stride, pixels += 3) {
-				pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale;
-				pixels[1] = max(0.0f, var[1] - mean[1]*mean[1]*invsample)*scale;
-				pixels[2] = max(0.0f, var[2] - mean[2]*mean[2]*invsample)*scale;
+				pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale_exposure;
+				pixels[1] = max(0.0f, var[1] - mean[1]*mean[1]*invsample)*scale_exposure;
+				pixels[2] = max(0.0f, var[2] - mean[2]*mean[2]*invsample)*scale_exposure;
 			}
 		}
 		else {
@@ -200,14 +201,28 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
 
 		if(components == 1) {
 			for(int i = 0; i < size; i++, in += pass_stride, pixels++) {
-				pixels[0] = in[0]*scale;
+				pixels[0] = in[0]*scale_exposure;
 			}
 		}
 		else if(components == 3) {
 			for(int i = 0; i < size; i++, in += pass_stride, pixels += 3) {
-				pixels[0] = in[0]*scale;
-				pixels[1] = in[1]*scale;
-				pixels[2] = in[2]*scale;
+				pixels[0] = in[0]*scale_exposure;
+				pixels[1] = in[1]*scale_exposure;
+				pixels[2] = in[2]*scale_exposure;
+			}
+		}
+		else if(components == 4) {
+			assert(offset == DENOISING_PASS_COLOR);
+
+			/* Since the alpha channel is not involved in denoising, output the Combined alpha channel. */
+			assert(params.passes[0].type == PASS_COMBINED);
+			float *in_combined = buffer.data();
+
+			for(int i = 0; i < size; i++, in += pass_stride, in_combined += pass_stride, pixels += 4) {
+				pixels[0] = in[0]*scale_exposure;
+				pixels[1] = in[1]*scale_exposure;
+				pixels[2] = in[2]*scale_exposure;
+				pixels[3] = saturate(in_combined[3]*scale);
 			}
 		}
 		else {
@@ -218,7 +233,7 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
 	return true;
 }
 
-bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels)
+bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels, const string &name)
 {
 	if(buffer.data() == NULL) {
 		return false;
@@ -234,6 +249,14 @@ bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int
 			continue;
 		}
 
+		/* Tell Cryptomatte passes apart by their name. */
+		if(pass.type == PASS_CRYPTOMATTE) {
+			if(pass.name != name) {
+				pass_offset += pass.components;
+				continue;
+			}
+		}
+
 		float *in = buffer.data() + pass_offset;
 		int pass_stride = params.get_passes_size();
 
@@ -370,6 +393,17 @@ bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int
 					pixels[3] = f.w*invw;
 				}
 			}
+			else if(type == PASS_CRYPTOMATTE) {
+				for(int i = 0; i < size; i++, in += pass_stride, pixels += 4) {
+					float4 f = make_float4(in[0], in[1], in[2], in[3]);
+					/* x and z contain integer IDs, don't rescale them.
+					   y and w contain matte weights, they get scaled. */
+					pixels[0] = f.x;
+					pixels[1] = f.y * scale;
+					pixels[2] = f.z;
+					pixels[3] = f.w * scale;
+				}
+			}
 			else {
 				for(int i = 0; i < size; i++, in += pass_stride, pixels += 4) {
 					float4 f = make_float4(in[0], in[1], in[2], in[3]);
diff --git a/intern/cycles/render/buffers.h b/intern/cycles/render/buffers.h
index 1b06ffe33a6..a8f019dddd6 100644
--- a/intern/cycles/render/buffers.h
+++ b/intern/cycles/render/buffers.h
@@ -50,7 +50,7 @@ public:
 	int full_height;
 
 	/* passes */
-	array<Pass> passes;
+	vector<Pass> passes;
 	bool denoising_data_pass;
 	/* If only some light path types should be denoised, an additional pass is needed. */
 	bool denoising_clean_pass;
@@ -84,7 +84,7 @@ public:
 	void zero();
 
 	bool copy_from_device();
-	bool get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels);
+	bool get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels, const string &name);
 	bool get_denoising_pass_rect(int offset, float exposure, int sample, int components, float *pixels);
 };
 
diff --git a/intern/cycles/render/coverage.cpp b/intern/cycles/render/coverage.cpp
new file mode 100644
index 00000000000..72ef4cda3ff
--- /dev/null
+++ b/intern/cycles/render/coverage.cpp
@@ -0,0 +1,143 @@
+/*
+ * Copyright 2018 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "render/coverage.h"
+#include "kernel/kernel_compat_cpu.h"
+#include "kernel/split/kernel_split_data.h"
+#include "kernel/kernel_globals.h"
+#include "kernel/kernel_id_passes.h"
+#include "kernel/kernel_types.h"
+#include "util/util_map.h"
+#include "util/util_vector.h"
+
+CCL_NAMESPACE_BEGIN
+
+static bool crypomatte_comp(const pair<float, float>& i, const pair<float, float> j) { return i.first > j.first; }
+
+void Coverage::finalize()
+{
+	int pass_offset = 0;
+	if(kernel_data.film.cryptomatte_passes & CRYPT_OBJECT) {
+		finalize_buffer(coverage_object, pass_offset);
+		pass_offset += kernel_data.film.cryptomatte_depth * 4;
+	}
+	if(kernel_data.film.cryptomatte_passes & CRYPT_MATERIAL) {
+		finalize_buffer(coverage_material, pass_offset);
+		pass_offset += kernel_data.film.cryptomatte_depth * 4;
+	}
+	if(kernel_data.film.cryptomatte_passes & CRYPT_ASSET) {
+		finalize_buffer(coverage_asset, pass_offset);
+	}
+}
+
+void Coverage::init_path_trace()
+{
+	kg->coverage_object = kg->coverage_material =  kg->coverage_asset = NULL;
+
+	if(kernel_data.film.cryptomatte_passes & CRYPT_ACCURATE) {
+		if(kernel_data.film.cryptomatte_passes & CRYPT_OBJECT) {
+			coverage_object.clear();
+			coverage_object.resize(tile.w * tile.h);
+		}
+		if(kernel_data.film.cryptomatte_passes & CRYPT_MATERIAL) {
+			coverage_material.clear();
+			coverage_material.resize(tile.w * tile.h);
+		}
+		if(kernel_data.film.cryptomatte_passes & CRYPT_ASSET) {
+			coverage_asset.clear();
+			coverage_asset.resize(tile.w * tile.h);
+		}
+	}
+}
+
+void Coverage::init_pixel(int x, int y)
+{
+	if(kernel_data.film.cryptomatte_passes & CRYPT_ACCURATE) {
+		const int pixel_index = tile.w * (y - tile.y) + x - tile.x;
+		if(kernel_data.film.cryptomatte_passes & CRYPT_OBJECT) {
+			kg->coverage_object = &coverage_object[pixel_index];
+		}
+		if(kernel_data.film.cryptomatte_passes & CRYPT_MATERIAL) {
+			kg->coverage_material = &coverage_material[pixel_index];
+		}
+		if(kernel_data.film.cryptomatte_passes & CRYPT_ASSET) {
+			kg->coverage_asset = &coverage_asset[pixel_index];
+		}
+	}
+}
+
+void Coverage::finalize_buffer(vector<CoverageMap> & coverage, const int pass_offset)
+{
+	if(kernel_data.film.cryptomatte_passes & CRYPT_ACCURATE) {
+		flatten_buffer(coverage, pass_offset);
+	}
+	else {
+		sort_buffer(pass_offset);
+	}
+}
+
+void Coverage::flatten_buffer(vector<CoverageMap> &coverage, const int pass_offset)
+{
+	/* Sort the coverage map and write it to the output */
+	int pixel_index = 0;
+	int pass_stride = tile.buffers->params.get_passes_size();
+	for(int y = 0; y < tile.h; ++y) {
+		for(int x = 0; x < tile.w; ++x) {
+			const CoverageMap& pixel = coverage[pixel_index];
+			if(!pixel.empty()) {
+				/* buffer offset */
+				int index = x + y * tile.stride;
+				float *buffer = (float*)tile.buffer + index*pass_stride;
+
+				/* sort the cryptomatte pixel */
+				vector<pair<float, float> > sorted_pixel;
+				for(CoverageMap::const_iterator it = pixel.begin(); it != pixel.end(); ++it) {
+					sorted_pixel.push_back(std::make_pair(it->second, it->first));
+				}
+				sort(sorted_pixel.begin(), sorted_pixel.end(), crypomatte_comp);
+				int num_slots = 2 * (kernel_data.film.cryptomatte_depth);
+				if(sorted_pixel.size() > num_slots) {
+					float leftover = 0.0f;
+					for(vector<pair<float, float> >::iterator it = sorted_pixel.begin()+num_slots; it != sorted_pixel.end(); ++it) {
+						leftover += it->first;
+					}
+					sorted_pixel[num_slots-1].first += leftover;
+				}
+				int limit = min(num_slots, sorted_pixel.size());
+				for(int i = 0; i < limit; ++i) {
+					kernel_write_id_slots(buffer + kernel_data.film.pass_cryptomatte + pass_offset, 2 * (kernel_data.film.cryptomatte_depth), sorted_pixel[i].second, sorted_pixel[i].first);
+				}
+			}
+			++pixel_index;
+		}
+	}
+}
+
+void Coverage::sort_buffer(const int pass_offset)
+{
+	/* Sort the coverage map and write it to the output */
+	int pass_stride = tile.buffers->params.get_passes_size();
+	for(int y = 0; y < tile.h; ++y) {
+		for(int x = 0; x < tile.w; ++x) {
+			/* buffer offset */
+			int index = x + y*tile.stride;
+			float *buffer = (float*)tile.buffer + index*pass_stride;
+			kernel_sort_id_slots(buffer + kernel_data.film.pass_cryptomatte + pass_offset, 2 * (kernel_data.film.cryptomatte_depth));
+		}
+	}
+}
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/render/coverage.h b/intern/cycles/render/coverage.h
new file mode 100644
index 00000000000..16176ce4beb
--- /dev/null
+++ b/intern/cycles/render/coverage.h
@@ -0,0 +1,49 @@
+/*
+ * Copyright 2018 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include "render/buffers.h"
+#include "kernel/kernel_compat_cpu.h"
+#include "kernel/split/kernel_split_data.h"
+#include "kernel/kernel_globals.h"
+#include "util/util_map.h"
+#include "util/util_vector.h"
+
+#ifndef __COVERAGE_H__
+#define __COVERAGE_H__
+
+CCL_NAMESPACE_BEGIN
+
+class Coverage {
+public:
+	Coverage(KernelGlobals *kg_, RenderTile &tile_) : kg(kg_), tile(tile_) { }
+	void init_path_trace();
+	void init_pixel(int x, int y);
+	void finalize();
+private:
+	vector<CoverageMap>coverage_object;
+	vector<CoverageMap>coverage_material;
+	vector<CoverageMap>coverage_asset;
+	KernelGlobals *kg;
+	RenderTile &tile;
+	void finalize_buffer(vector<CoverageMap>&coverage, const int pass_offset);
+	void flatten_buffer(vector<CoverageMap>&coverage, const int pass_offset);
+	void sort_buffer(const int pass_offset);
+};
+
+
+CCL_NAMESPACE_END
+
+#endif /* __COVERAGE_H__ */
diff --git a/intern/cycles/render/film.cpp b/intern/cycles/render/film.cpp
index 8f3596ade58..d0f15496e50 100644
--- a/intern/cycles/render/film.cpp
+++ b/intern/cycles/render/film.cpp
@@ -38,11 +38,14 @@ static bool compare_pass_order(const Pass& a, const Pass& b)
 	return (a.components > b.components);
 }
 
-void Pass::add(PassType type, array<Pass>& passes)
+void Pass::add(PassType type, vector<Pass>& passes, const char *name)
 {
-	for(size_t i = 0; i < passes.size(); i++)
-		if(passes[i].type == type)
+	for(size_t i = 0; i < passes.size(); i++) {
+		if(passes[i].type == type &&
+		   (name ? (passes[i].name == name) : passes[i].name.empty())) {
 			return;
+		}
+	}
 
 	Pass pass;
 
@@ -50,6 +53,9 @@ void Pass::add(PassType type, array<Pass>& passes)
 	pass.filter = true;
 	pass.exposure = false;
 	pass.divide_type = PASS_NONE;
+	if(name) {
+		pass.name = name;
+	}
 
 	switch(type) {
 		case PASS_NONE:
@@ -155,13 +161,15 @@ void Pass::add(PassType type, array<Pass>& passes)
 			pass.components = 4;
 			pass.exposure = true;
 			break;
-
+		case PASS_CRYPTOMATTE:
+			pass.components = 4;
+			break;
 		default:
 			assert(false);
 			break;
 	}
 
-	passes.push_back_slow(pass);
+	passes.push_back(pass);
 
 	/* order from by components, to ensure alignment so passes with size 4
 	 * come first and then passes with size 1 */
@@ -171,19 +179,19 @@ void Pass::add(PassType type, array<Pass>& passes)
 		Pass::add(pass.divide_type, passes);
 }
 
-bool Pass::equals(const array<Pass>& A, const array<Pass>& B)
+bool Pass::equals(const vector<Pass>& A, const vector<Pass>& B)
 {
 	if(A.size() != B.size())
 		return false;
 
 	for(int i = 0; i < A.size(); i++)
-		if(A[i].type != B[i].type)
+		if(A[i].type != B[i].type || A[i].name != B[i].name)
 			return false;
 
 	return true;
 }
 
-bool Pass::contains(const array<Pass>& passes, PassType type)
+bool Pass::contains(const vector<Pass>& passes, PassType type)
 {
 	for(size_t i = 0; i < passes.size(); i++)
 		if(passes[i].type == type)
@@ -290,6 +298,7 @@ Film::Film()
 
 	use_light_visibility = false;
 	filter_table_offset = TABLE_OFFSET_INVALID;
+	cryptomatte_passes = CRYPT_NONE;
 
 	need_update = true;
 }
@@ -314,6 +323,8 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
 	kfilm->pass_stride = 0;
 	kfilm->use_light_pass = use_light_visibility || use_sample_clamp;
 
+	bool have_cryptomatte = false;
+
 	for(size_t i = 0; i < passes.size(); i++) {
 		Pass& pass = passes[i];
 
@@ -434,7 +445,10 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
 #endif
 			case PASS_RENDER_TIME:
 				break;
-
+			case PASS_CRYPTOMATTE:
+				kfilm->pass_cryptomatte = have_cryptomatte ? min(kfilm->pass_cryptomatte, kfilm->pass_stride) : kfilm->pass_stride;
+				have_cryptomatte = true;
+				break;
 			default:
 				assert(false);
 				break;
@@ -471,6 +485,9 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
 	kfilm->mist_inv_depth = (mist_depth > 0.0f)? 1.0f/mist_depth: 0.0f;
 	kfilm->mist_falloff = mist_falloff;
 
+	kfilm->cryptomatte_passes = cryptomatte_passes;
+	kfilm->cryptomatte_depth = cryptomatte_depth;
+
 	pass_stride = kfilm->pass_stride;
 	denoising_data_offset = kfilm->pass_denoising_data;
 	denoising_clean_offset = kfilm->pass_denoising_clean;
@@ -490,7 +507,7 @@ bool Film::modified(const Film& film)
 	return !Node::equals(film) || !Pass::equals(passes, film.passes);
 }
 
-void Film::tag_passes_update(Scene *scene, const array<Pass>& passes_)
+void Film::tag_passes_update(Scene *scene, const vector<Pass>& passes_)
 {
 	if(Pass::contains(passes, PASS_UV) != Pass::contains(passes_, PASS_UV)) {
 		scene->mesh_manager->tag_update(scene);
diff --git a/intern/cycles/render/film.h b/intern/cycles/render/film.h
index 6ab2eea79b8..57f1bf4eb64 100644
--- a/intern/cycles/render/film.h
+++ b/intern/cycles/render/film.h
@@ -45,10 +45,11 @@ public:
 	bool filter;
 	bool exposure;
 	PassType divide_type;
+	string name;
 
-	static void add(PassType type, array<Pass>& passes);
-	static bool equals(const array<Pass>& A, const array<Pass>& B);
-	static bool contains(const array<Pass>& passes, PassType);
+	static void add(PassType type, vector<Pass>& passes, const char* name = NULL);
+	static bool equals(const vector<Pass>& A, const vector<Pass>& B);
+	static bool contains(const vector<Pass>& passes, PassType);
 };
 
 class Film : public Node {
@@ -56,7 +57,7 @@ public:
 	NODE_DECLARE
 
 	float exposure;
-	array<Pass> passes;
+	vector<Pass> passes;
 	bool denoising_data_pass;
 	bool denoising_clean_pass;
 	int denoising_flags;
@@ -76,6 +77,8 @@ public:
 
 	bool use_light_visibility;
 	bool use_sample_clamp;
+	CryptomatteType cryptomatte_passes;
+	int cryptomatte_depth;
 
 	bool need_update;
 
@@ -86,7 +89,7 @@ public:
 	void device_free(Device *device, DeviceScene *dscene, Scene *scene);
 
 	bool modified(const Film& film);
-	void tag_passes_update(Scene *scene, const array<Pass>& passes_);
+	void tag_passes_update(Scene *scene, const vector<Pass>& passes_);
 	void tag_update(Scene *scene);
 };
 
diff --git a/intern/cycles/render/object.cpp b/intern/cycles/render/object.cpp
index e3f35c366d6..a56a8a6ec58 100644
--- a/intern/cycles/render/object.cpp
+++ b/intern/cycles/render/object.cpp
@@ -28,6 +28,7 @@
 #include "util/util_map.h"
 #include "util/util_progress.h"
 #include "util/util_vector.h"
+#include "util/util_murmurhash.h"
 
 #include "subd/subd_patch_table.h"
 
@@ -483,6 +484,10 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
 	kobject.numverts = mesh->verts.size();
 	kobject.patch_map_offset = 0;
 	kobject.attribute_map_offset = 0;
+	uint32_t hash_name = util_murmur_hash3(ob->name.c_str(), ob->name.length(), 0);
+	uint32_t hash_asset = util_murmur_hash3(ob->asset_name.c_str(), ob->asset_name.length(), 0);
+	kobject.cryptomatte_object = util_hash_to_float(hash_name);
+	kobject.cryptomatte_asset = util_hash_to_float(hash_asset);
 
 	/* Object flag. */
 	if(ob->use_holdout) {
diff --git a/intern/cycles/render/object.h b/intern/cycles/render/object.h
index b80c4aef70b..bd44b35aba3 100644
--- a/intern/cycles/render/object.h
+++ b/intern/cycles/render/object.h
@@ -48,6 +48,7 @@ public:
 	BoundBox bounds;
 	uint random_id;
 	int pass_id;
+	ustring asset_name;
 	vector<ParamValue> attributes;
 	uint visibility;
 	array<Transform> motion;
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index d0aa985b035..d6ecafa19b7 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -682,7 +682,10 @@ DeviceRequestedFeatures Session::get_requested_device_features()
 	BakeManager *bake_manager = scene->bake_manager;
 	requested_features.use_baking = bake_manager->get_baking();
 	requested_features.use_integrator_branched = (scene->integrator->method == Integrator::BRANCHED_PATH);
-	requested_features.use_denoising = params.use_denoising;
+	if(params.denoising_passes) {
+		requested_features.use_denoising = true;
+		requested_features.use_shadow_tricks = true;
+	}
 
 	return requested_features;
 }
diff --git a/intern/cycles/render/session.h b/intern/cycles/render/session.h
index 61f62f8e712..d2ec13ada1f 100644
--- a/intern/cycles/render/session.h
+++ b/intern/cycles/render/session.h
@@ -58,6 +58,7 @@ public:
 	bool display_buffer_linear;
 
 	bool use_denoising;
+	bool denoising_passes;
 	int denoising_radius;
 	float denoising_strength;
 	float denoising_feature_strength;
@@ -89,6 +90,7 @@ public:
 		threads = 0;
 
 		use_denoising = false;
+		denoising_passes = false;
 		denoising_radius = 8;
 		denoising_strength = 0.0f;
 		denoising_feature_strength = 0.0f;
diff --git a/intern/cycles/render/shader.cpp b/intern/cycles/render/shader.cpp
index ac605305b94..8d0cec7b14e 100644
--- a/intern/cycles/render/shader.cpp
+++ b/intern/cycles/render/shader.cpp
@@ -30,6 +30,7 @@
 #include "render/tables.h"
 
 #include "util/util_foreach.h"
+#include "util/util_murmurhash.h"
 
 #ifdef WITH_OCIO
 #  include <OpenColorIO/OpenColorIO.h>
@@ -523,12 +524,15 @@ void ShaderManager::device_update_common(Device *device,
 		if(shader->is_constant_emission(&constant_emission))
 			flag |= SD_HAS_CONSTANT_EMISSION;
 
+		uint32_t cryptomatte_id = util_murmur_hash3(shader->name.c_str(), shader->name.length(), 0);
+		
 		/* regular shader */
 		kshader->flags = flag;
 		kshader->pass_id = shader->pass_id;
 		kshader->constant_emission[0] = constant_emission.x;
 		kshader->constant_emission[1] = constant_emission.y;
 		kshader->constant_emission[2] = constant_emission.z;
+		kshader->cryptomatte_id = util_hash_to_float(cryptomatte_id);
 		kshader++;
 
 		has_transparent_shadow |= (flag & SD_HAS_TRANSPARENT_SHADOW) != 0;
diff --git a/intern/cycles/util/CMakeLists.txt b/intern/cycles/util/CMakeLists.txt
index 291f9a9fcae..4f623c5dfb7 100644
--- a/intern/cycles/util/CMakeLists.txt
+++ b/intern/cycles/util/CMakeLists.txt
@@ -15,6 +15,7 @@ set(SRC
 	util_logging.cpp
 	util_math_cdf.cpp
 	util_md5.cpp
+	util_murmurhash.cpp
 	util_path.cpp
 	util_string.cpp
 	util_simd.cpp
@@ -64,6 +65,7 @@ set(SRC_HEADERS
 	util_math_int4.h
 	util_math_matrix.h
 	util_md5.h
+	util_murmurhash.h
 	util_opengl.h
 	util_optimization.h
 	util_param.h
diff --git a/intern/cycles/util/util_atomic.h b/intern/cycles/util/util_atomic.h
index f3c7ae546a0..e17e99d0acd 100644
--- a/intern/cycles/util/util_atomic.h
+++ b/intern/cycles/util/util_atomic.h
@@ -23,6 +23,7 @@
 #include "atomic_ops.h"
 
 #define atomic_add_and_fetch_float(p, x) atomic_add_and_fetch_fl((p), (x))
+#define atomic_compare_and_swap_float(p, old_val, new_val) atomic_cas_float((p), (old_val), (new_val))
 
 #define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1)
 #define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_add_uint32((p), -1)
@@ -57,6 +58,20 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so
 	return new_value.float_value;
 }
 
+ccl_device_inline float atomic_compare_and_swap_float(volatile ccl_global float *dest,
+                                                      const float old_val, const float new_val)
+{
+	union {
+		unsigned int int_value;
+		float float_value;
+	} new_value, prev_value, result;
+	prev_value.float_value = old_val;
+	new_value.float_value = new_val;
+	result.int_value = atomic_cmpxchg((volatile ccl_global unsigned int *)dest,
+                                       prev_value.int_value, new_value.int_value);
+	return result.float_value;
+}
+
 #define atomic_fetch_and_add_uint32(p, x) atomic_add((p), (x))
 #define atomic_fetch_and_inc_uint32(p) atomic_inc((p))
 #define atomic_fetch_and_dec_uint32(p) atomic_dec((p))
@@ -75,6 +90,19 @@ ccl_device_inline float atomic_add_and_fetch_float(volatile ccl_global float *so
 #define atomic_fetch_and_inc_uint32(p) atomic_fetch_and_add_uint32((p), 1)
 #define atomic_fetch_and_dec_uint32(p) atomic_fetch_and_sub_uint32((p), 1)
 
+ccl_device_inline float atomic_compare_and_swap_float(volatile float *dest,
+                                                      const float old_val, const float new_val)
+{
+	union {
+		unsigned int int_value;
+		float float_value;
+	} new_value, prev_value, result;
+	prev_value.float_value = old_val;
+	new_value.float_value = new_val;
+	result.int_value = atomicCAS((unsigned int *)dest, prev_value.int_value,new_value.int_value);
+	return result.float_value;
+}
+
 #define CCL_LOCAL_MEM_FENCE
 #define ccl_barrier(flags) __syncthreads()
 
diff --git a/intern/cycles/util/util_murmurhash.cpp b/intern/cycles/util/util_murmurhash.cpp
new file mode 100644
index 00000000000..68b2f2031be
--- /dev/null
+++ b/intern/cycles/util/util_murmurhash.cpp
@@ -0,0 +1,127 @@
+/*
+ * Copyright 2018 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/* This is taken from alShaders/Cryptomatte/MurmurHash3.h:
+ *
+ * MurmurHash3 was written by Austin Appleby, and is placed in the public
+ * domain. The author hereby disclaims copyright to this source code.
+ *
+ */
+
+#include <stdlib.h>
+#include <string.h>
+
+#include "util/util_algorithm.h"
+#include "util/util_murmurhash.h"
+
+#if defined(_MSC_VER)
+#  define ROTL32(x,y) _rotl(x,y)
+#  define ROTL64(x,y) _rotl64(x,y)
+#  define BIG_CONSTANT(x) (x)
+#else
+ccl_device_inline uint32_t rotl32(uint32_t x, int8_t r)
+{
+	return (x << r) | (x >> (32 - r));
+}
+#  define ROTL32(x,y) rotl32(x,y)
+#  define BIG_CONSTANT(x) (x##LLU)
+#endif
+
+CCL_NAMESPACE_BEGIN
+
+/* Block read - if your platform needs to do endian-swapping or can only
+ * handle aligned reads, do the conversion here. */
+ccl_device_inline uint32_t mm_hash_getblock32(const uint32_t *p, int i)
+{
+	return p[i];
+}
+
+/* Finalization mix - force all bits of a hash block to avalanche */
+ccl_device_inline uint32_t mm_hash_fmix32 ( uint32_t h )
+{
+	h ^= h >> 16;
+	h *= 0x85ebca6b;
+	h ^= h >> 13;
+	h *= 0xc2b2ae35;
+	h ^= h >> 16;
+	return h;
+}
+
+uint32_t util_murmur_hash3(const void *key, int len, uint32_t seed)
+{
+	const uint8_t * data = (const uint8_t*)key;
+	const int nblocks = len / 4;
+
+	uint32_t h1 = seed;
+
+	const uint32_t c1 = 0xcc9e2d51;
+	const uint32_t c2 = 0x1b873593;
+
+	const uint32_t * blocks = (const uint32_t *)(data + nblocks*4);
+
+	for(int i = -nblocks; i; i++) {
+		uint32_t k1 = mm_hash_getblock32(blocks,i);
+
+		k1 *= c1;
+		k1 = ROTL32(k1,15);
+		k1 *= c2;
+
+		h1 ^= k1;
+		h1 = ROTL32(h1,13);
+		h1 = h1 * 5 + 0xe6546b64;
+	}
+
+	const uint8_t *tail = (const uint8_t*)(data + nblocks*4);
+
+	uint32_t k1 = 0;
+
+	switch(len & 3) {
+		case 3:
+			k1 ^= tail[2] << 16;
+			ATTR_FALLTHROUGH;
+		case 2:
+			k1 ^= tail[1] << 8;
+			ATTR_FALLTHROUGH;
+		case 1:
+			k1 ^= tail[0];
+			k1 *= c1;
+			k1 = ROTL32(k1,15);
+			k1 *= c2;
+			h1 ^= k1;
+	}
+
+	h1 ^= len;
+	h1 = mm_hash_fmix32(h1);
+	return h1;
+}
+
+/* This is taken from the cryptomatte specification 1.0 */
+float util_hash_to_float(uint32_t hash)
+{
+	uint32_t mantissa = hash & (( 1 << 23) - 1);
+	uint32_t exponent = (hash >> 23) & ((1 << 8) - 1);
+	exponent = max(exponent, (uint32_t) 1);
+	exponent = min(exponent, (uint32_t) 254);
+	exponent = exponent << 23;
+	uint32_t sign = (hash >> 31);
+	sign = sign << 31;
+	uint32_t float_bits = sign | exponent | mantissa;
+	float f;
+	memcpy(&f, &float_bits, sizeof(uint32_t));
+	return f;
+}
+
+CCL_NAMESPACE_END
diff --git a/intern/cycles/util/util_murmurhash.h b/intern/cycles/util/util_murmurhash.h
new file mode 100644
index 00000000000..824ed59cb16
--- /dev/null
+++ b/intern/cycles/util/util_murmurhash.h
@@ -0,0 +1,30 @@
+/*
+ * Copyright 2018 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+
+#ifndef __UTIL_MURMURHASH_H__
+#define __UTIL_MURMURHASH_H__
+
+#include "util/util_types.h"
+
+CCL_NAMESPACE_BEGIN
+
+uint32_t util_murmur_hash3(const void *key, int len, uint32_t seed);
+float util_hash_to_float(uint32_t hash);
+
+CCL_NAMESPACE_END
+
+#endif /* __UTIL_MURMURHASH_H__ */
diff --git a/intern/locale/boost_locale_wrapper.cpp b/intern/locale/boost_locale_wrapper.cpp
index 0707c0dd3e3..3fd8f146aa3 100644
--- a/intern/locale/boost_locale_wrapper.cpp
+++ b/intern/locale/boost_locale_wrapper.cpp
@@ -112,13 +112,17 @@ const char *bl_locale_pgettext(const char *msgctxt, const char *msgid)
 			return r;
 		return msgid;
 	}
-	catch(std::bad_cast const &e) { /* if std::has_facet<char_message_facet>(l) == false, LC_ALL = "C" case */
-//		std::cout << "bl_locale_pgettext(" << msgid << "): " << e.what() << " \n";
+	catch(const std::bad_cast &e) { /* if std::has_facet<char_message_facet>(l) == false, LC_ALL = "C" case */
+#ifndef NDEBUG
+		std::cout << "bl_locale_pgettext(" << msgid << "): " << e.what() << " \n";
+#endif
 		(void)e;
 		return msgid;
 	}
-	catch(std::exception const &e) {
-//		std::cout << "bl_locale_pgettext(" << msgctxt << ", " << msgid << "): " << e.what() << " \n";
+	catch(const std::exception &e) {
+#ifndef NDEBUG
+		std::cout << "bl_locale_pgettext(" << msgctxt << ", " << msgid << "): " << e.what() << " \n";
+#endif
 		(void)e;
 		return msgid;
 	}
diff --git a/release/scripts/modules/addon_utils.py b/release/scripts/modules/addon_utils.py
index dc7f4053a01..713a11362c6 100644
--- a/release/scripts/modules/addon_utils.py
+++ b/release/scripts/modules/addon_utils.py
@@ -472,7 +472,12 @@ def reset_all(*, reload_scripts=False):
 
 def disable_all():
     import sys
-    for mod_name, mod in sys.modules.items():
+    # Collect modules to disable first because dict can be modified as we disable.
+    addon_modules = [
+        item for item in sys.modules.items()
+        if getattr(item[1], "__addon_enabled__", False)
+    ]
+    for mod_name, mod in addon_modules:
         if getattr(mod, "__addon_enabled__", False):
             disable(mod_name)
 
diff --git a/release/scripts/startup/bl_operators/uvcalc_lightmap.py b/release/scripts/startup/bl_operators/uvcalc_lightmap.py
index d1774205aee..11eb7e1eb88 100644
--- a/release/scripts/startup/bl_operators/uvcalc_lightmap.py
+++ b/release/scripts/startup/bl_operators/uvcalc_lightmap.py
@@ -77,9 +77,9 @@ class prettyface:
             # f, (len_min, len_mid, len_max)
             self.uv = data
 
-            f1, lens1, lens1ord = data[0]
+            _f1, lens1, lens1ord = data[0]
             if data[1]:
-                f2, lens2, lens2ord = data[1]
+                _f2, lens2, lens2ord = data[1]
                 self.width = (lens1[lens1ord[0]] + lens2[lens2ord[0]]) / 2.0
                 self.height = (lens1[lens1ord[1]] + lens2[lens2ord[1]]) / 2.0
             else:  # 1 tri :/
@@ -205,12 +205,12 @@ class prettyface:
                     fuv[I[0]][:] = p2
                     fuv[I[1]][:] = p3
 
-            f, lens, lensord = uv[0]
+            f = uv[0][0]
 
             set_uv(f, (x1, y1), (x1, y2 - margin_h), (x2 - margin_w, y1))
 
             if uv[1]:
-                f, lens, lensord = uv[1]
+                f = uv[1][0]
                 set_uv(f, (x2, y2), (x2, y1 + margin_h), (x1 + margin_w, y2))
 
         else:  # 1 QUAD
diff --git a/release/scripts/startup/bl_operators/uvcalc_smart_project.py b/release/scripts/startup/bl_operators/uvcalc_smart_project.py
index 16a24736398..ec6a9fc92e3 100644
--- a/release/scripts/startup/bl_operators/uvcalc_smart_project.py
+++ b/release/scripts/startup/bl_operators/uvcalc_smart_project.py
@@ -143,9 +143,9 @@ def island2Edge(island):
     unique_points = {}
 
     for f in island:
-        f_uvkey = map(tuple, f.uv)
+        f_uvkey = list(map(tuple, f.uv))
 
-        for vIdx, edkey in enumerate(f.edge_keys):
+        for vIdx in range(len(f_uvkey)):
             unique_points[f_uvkey[vIdx]] = f.uv[vIdx]
 
             if f.v[vIdx].index > f.v[vIdx - 1].index:
@@ -158,18 +158,14 @@ def island2Edge(island):
             try:
                 edges[f_uvkey[i1], f_uvkey[i2]] *= 0  # sets any edge with more than 1 user to 0 are not returned.
             except:
-                edges[f_uvkey[i1], f_uvkey[i2]] = (f.uv[i1] - f.uv[i2]).length,
+                edges[f_uvkey[i1], f_uvkey[i2]] = (f.uv[i1] - f.uv[i2]).length
 
     # If 2 are the same then they will be together, but full [a,b] order is not correct.
 
     # Sort by length
-
     length_sorted_edges = [(Vector(key[0]), Vector(key[1]), value) for key, value in edges.items() if value != 0]
 
-    try:
-        length_sorted_edges.sort(key=lambda A: -A[2])  # largest first
-    except:
-        length_sorted_edges.sort(lambda A, B: cmp(B[2], A[2]))
+    length_sorted_edges.sort(key=lambda a: -a[2])  # largest first
 
     # Its okay to leave the length in there.
     # for e in length_sorted_edges:
diff --git a/release/scripts/startup/bl_ui/properties_data_modifier.py b/release/scripts/startup/bl_ui/properties_data_modifier.py
index 1a80a160ea5..23ab644cba1 100644
--- a/release/scripts/startup/bl_ui/properties_data_modifier.py
+++ b/release/scripts/startup/bl_ui/properties_data_modifier.py
@@ -523,6 +523,9 @@ class DATA_PT_modifiers(ModifierButtonsPanel, Panel):
             sub.active = bool(md.vertex_group)
             sub.prop(md, "invert_vertex_group", text="", icon='ARROW_LEFTRIGHT')
 
+        col = layout.column()
+        col.prop(md, "threshold")
+
     def MESH_DEFORM(self, layout, ob, md):
         split = layout.split()
 
diff --git a/source/blender/blenkernel/BKE_global.h b/source/blender/blenkernel/BKE_global.h
index 0ce560532d5..6ed32fc43ac 100644
--- a/source/blender/blenkernel/BKE_global.h
+++ b/source/blender/blenkernel/BKE_global.h
@@ -121,14 +121,14 @@ typedef struct Global {
 
 /* G.debug */
 enum {
-	G_DEBUG =           (1 << 0), /* general debug flag, print more info in unexpected cases */
+	G_DEBUG =           (1 << 0),  /* general debug flag, print more info in unexpected cases */
 	G_DEBUG_FFMPEG =    (1 << 1),
-	G_DEBUG_PYTHON =    (1 << 2), /* extra python info */
-	G_DEBUG_EVENTS =    (1 << 3), /* input/window/screen events */
-	G_DEBUG_HANDLERS =  (1 << 4), /* events handling */
-	G_DEBUG_WM =        (1 << 5), /* operator, undo */
-	G_DEBUG_JOBS =      (1 << 6), /* jobs time profiling */
-	G_DEBUG_FREESTYLE = (1 << 7), /* freestyle messages */
+	G_DEBUG_PYTHON =    (1 << 2),  /* extra python info */
+	G_DEBUG_EVENTS =    (1 << 3),  /* input/window/screen events */
+	G_DEBUG_HANDLERS =  (1 << 4),  /* events handling */
+	G_DEBUG_WM =        (1 << 5),  /* operator, undo */
+	G_DEBUG_JOBS =      (1 << 6),  /* jobs time profiling */
+	G_DEBUG_FREESTYLE = (1 << 7),  /* freestyle messages */
 	G_DEBUG_DEPSGRAPH_BUILD      = (1 << 8),   /* depsgraph construction messages */
 	G_DEBUG_DEPSGRAPH_EVAL       = (1 << 9),   /* depsgraph evaluation messages */
 	G_DEBUG_DEPSGRAPH_TAG        = (1 << 10),  /* depsgraph tagging messages */
@@ -139,11 +139,11 @@ enum {
 	                     G_DEBUG_DEPSGRAPH_EVAL |
 	                     G_DEBUG_DEPSGRAPH_TAG |
 	                     G_DEBUG_DEPSGRAPH_TIME),
-	G_DEBUG_SIMDATA =   (1 << 14), /* sim debug data display */
-	G_DEBUG_GPU_MEM =   (1 << 15), /* gpu memory in status bar */
-	G_DEBUG_GPU =       (1 << 16), /* gpu debug */
-	G_DEBUG_IO = (1 << 17),   /* IO Debugging (for Collada, ...)*/
-	G_DEBUG_GPU_SHADERS = (1 << 18),   /* GLSL shaders */
+	G_DEBUG_SIMDATA =   (1 << 14),  /* sim debug data display */
+	G_DEBUG_GPU_MEM =   (1 << 15),  /* gpu memory in status bar */
+	G_DEBUG_GPU =       (1 << 16),  /* gpu debug */
+	G_DEBUG_IO =        (1 << 17),  /* IO Debugging (for Collada, ...)*/
+	G_DEBUG_GPU_SHADERS = (1 << 18),  /* GLSL shaders */
 };
 
 #define G_DEBUG_ALL  (G_DEBUG | G_DEBUG_FFMPEG | G_DEBUG_PYTHON | G_DEBUG_EVENTS | G_DEBUG_WM | G_DEBUG_JOBS | \
diff --git a/source/blender/blenkernel/intern/blendfile.c b/source/blender/blenkernel/intern/blendfile.c
index aa5530704c6..c92648da67c 100644
--- a/source/blender/blenkernel/intern/blendfile.c
+++ b/source/blender/blenkernel/intern/blendfile.c
@@ -590,10 +590,6 @@ bool BKE_blendfile_write_partial(
 	 * (otherwise main->name will not be set at read time). */
 	BLI_strncpy(bmain_dst->name, bmain_src->name, sizeof(bmain_dst->name));
 
-	if (write_flags & G_FILE_RELATIVE_REMAP) {
-		path_list_backup = BKE_bpath_list_backup(bmain_src, path_list_flag);
-	}
-
 	BLO_main_expander(blendfile_write_partial_cb);
 	BLO_expand_main(NULL, bmain_src);
 
@@ -613,10 +609,27 @@ bool BKE_blendfile_write_partial(
 		}
 	}
 
+	/* Backup paths because remap relative will overwrite them.
+	 *
+	 * NOTE: we do this only on the list of datablocks that we are writing
+	 * because the restored full list is not guaranteed to be in the same
+	 * order as before, as expected by BKE_bpath_list_restore.
+	 *
+	 * This happens because id_sort_by_name does not take into account
+	 * string case or the library name, so the order is not strictly
+	 * defined for two linked datablocks with the same name! */
+	if (write_flags & G_FILE_RELATIVE_REMAP) {
+		path_list_backup = BKE_bpath_list_backup(bmain_dst, path_list_flag);
+	}
 
 	/* save the buffer */
 	retval = BLO_write_file(bmain_dst, filepath, write_flags, reports, NULL);
 
+	if (path_list_backup) {
+		BKE_bpath_list_restore(bmain_dst, path_list_flag, path_list_backup);
+		BKE_bpath_list_free(path_list_backup);
+	}
+
 	/* move back the main, now sorted again */
 	set_listbasepointers(bmain_src, lbarray_dst);
 	a = set_listbasepointers(bmain_dst, lbarray_src);
@@ -632,11 +645,6 @@ bool BKE_blendfile_write_partial(
 
 	MEM_freeN(bmain_dst);
 
-	if (path_list_backup) {
-		BKE_bpath_list_restore(bmain_src, path_list_flag, path_list_backup);
-		BKE_bpath_list_free(path_list_backup);
-	}
-
 	return retval;
 }
 
diff --git a/source/blender/blenkernel/intern/scene.c b/source/blender/blenkernel/intern/scene.c
index 1eb65519596..3a067221aad 100644
--- a/source/blender/blenkernel/intern/scene.c
+++ b/source/blender/blenkernel/intern/scene.c
@@ -402,6 +402,7 @@ Scene *BKE_scene_copy(Main *bmain, Scene *sce, int type)
 			for (SceneRenderLayer *srl_dst = sce_copy->r.layers.first; srl_dst; srl_dst = srl_dst->next) {
 				for (FreestyleLineSet *lineset = srl_dst->freestyleConfig.linesets.first; lineset; lineset = lineset->next) {
 					if (lineset->linestyle) {
+						id_us_min(&lineset->linestyle->id);
 						/* XXX Not copying anim/actions here? */
 						BKE_id_copy_ex(bmain, (ID *)lineset->linestyle, (ID **)&lineset->linestyle, 0, false);
 					}
@@ -410,12 +411,14 @@ Scene *BKE_scene_copy(Main *bmain, Scene *sce, int type)
 
 			/* Full copy of world (included animations) */
 			if (sce_copy->world) {
+				id_us_min(&sce_copy->world->id);
 				BKE_id_copy_ex(bmain, (ID *)sce_copy->world, (ID **)&sce_copy->world, LIB_ID_COPY_ACTIONS, false);
 			}
 
 			/* Full copy of GreasePencil. */
 			/* XXX Not copying anim/actions here? */
 			if (sce_copy->gpd) {
+				id_us_min(&sce_copy->gpd->id);
 				BKE_id_copy_ex(bmain, (ID *)sce_copy->gpd, (ID **)&sce_copy->gpd, 0, false);
 			}
 		}
@@ -430,6 +433,7 @@ Scene *BKE_scene_copy(Main *bmain, Scene *sce, int type)
 		 * are done outside of blenkernel with ED_objects_single_users! */
 
 		/*  camera */
+		/* XXX This is most certainly useless? Object have not yet been duplicated... */
 		if (ELEM(type, SCE_COPY_LINK_DATA, SCE_COPY_FULL)) {
 			ID_NEW_REMAP(sce_copy->camera);
 		}
diff --git a/source/blender/blenloader/intern/readfile.c b/source/blender/blenloader/intern/readfile.c
index 08a45f367f4..9478c2e0351 100644
--- a/source/blender/blenloader/intern/readfile.c
+++ b/source/blender/blenloader/intern/readfile.c
@@ -5683,7 +5683,13 @@ static void direct_link_object(FileData *fd, Object *ob)
 	CLAMP(ob->rotmode, ROT_MODE_MIN, ROT_MODE_MAX);
 
 	if (ob->sculpt) {
-		ob->sculpt = MEM_callocN(sizeof(SculptSession), "reload sculpt session");
+		if (ob->mode & OB_MODE_ALL_SCULPT) {
+			ob->sculpt = MEM_callocN(sizeof(SculptSession), "reload sculpt session");
+			ob->sculpt->mode_type = ob->mode;
+		}
+		else {
+			ob->sculpt = NULL;
+		}
 	}
 
 	link_list(fd, &ob->lodlevels);
diff --git a/source/blender/bmesh/intern/bmesh_mesh.c b/source/blender/bmesh/intern/bmesh_mesh.c
index 185e5412b3d..4568c3b7965 100644
--- a/source/blender/bmesh/intern/bmesh_mesh.c
+++ b/source/blender/bmesh/intern/bmesh_mesh.c
@@ -1455,28 +1455,6 @@ void BM_mesh_elem_table_free(BMesh *bm, const char htype)
 	}
 }
 
-BMVert *BM_vert_at_index(BMesh *bm, const int index)
-{
-	BLI_assert((index >= 0) && (index < bm->totvert));
-	BLI_assert((bm->elem_table_dirty & BM_VERT) == 0);
-	return bm->vtable[index];
-}
-
-BMEdge *BM_edge_at_index(BMesh *bm, const int index)
-{
-	BLI_assert((index >= 0) && (index < bm->totedge));
-	BLI_assert((bm->elem_table_dirty & BM_EDGE) == 0);
-	return bm->etable[index];
-}
-
-BMFace *BM_face_at_index(BMesh *bm, const int index)
-{
-	BLI_assert((index >= 0) && (index < bm->totface));
-	BLI_assert((bm->elem_table_dirty & BM_FACE) == 0);
-	return bm->ftable[index];
-}
-
-
 BMVert *BM_vert_at_index_find(BMesh *bm, const int index)
 {
 	return BLI_mempool_findelem(bm->vpool, index);
diff --git a/source/blender/bmesh/intern/bmesh_mesh.h b/source/blender/bmesh/intern/bmesh_mesh.h
index d449aac04f5..9d7866c280a 100644
--- a/source/blender/bmesh/intern/bmesh_mesh.h
+++ b/source/blender/bmesh/intern/bmesh_mesh.h
@@ -75,9 +75,24 @@ void           BM_mesh_elem_table_ensure(BMesh *bm, const char htype);
 void           BM_mesh_elem_table_init(BMesh *bm, const char htype);
 void           BM_mesh_elem_table_free(BMesh *bm, const char htype);
 
-BMVert *BM_vert_at_index(BMesh *bm, const int index);
-BMEdge *BM_edge_at_index(BMesh *bm, const int index);
-BMFace *BM_face_at_index(BMesh *bm, const int index);
+BLI_INLINE BMVert *BM_vert_at_index(BMesh *bm, const int index)
+{
+	BLI_assert((index >= 0) && (index < bm->totvert));
+	BLI_assert((bm->elem_table_dirty & BM_VERT) == 0);
+	return bm->vtable[index];
+}
+BLI_INLINE BMEdge *BM_edge_at_index(BMesh *bm, const int index)
+{
+	BLI_assert((index >= 0) && (index < bm->totedge));
+	BLI_assert((bm->elem_table_dirty & BM_EDGE) == 0);
+	return bm->etable[index];
+}
+BLI_INLINE BMFace *BM_face_at_index(BMesh *bm, const int index)
+{
+	BLI_assert((index >= 0) && (index < bm->totface));
+	BLI_assert((bm->elem_table_dirty & BM_FACE) == 0);
+	return bm->ftable[index];
+}
 
 BMVert *BM_vert_at_index_find(BMesh *bm, const int index);
 BMEdge *BM_edge_at_index_find(BMesh *bm, const int index);
diff --git a/source/blender/editors/space_file/fsmenu.c b/source/blender/editors/space_file/fsmenu.c
index 38b8dd37d46..f925409d882 100644
--- a/source/blender/editors/space_file/fsmenu.c
+++ b/source/blender/editors/space_file/fsmenu.c
@@ -605,6 +605,8 @@ void fsmenu_read_system(struct FSMenu *fsmenu, int read_bookmarks)
 					/* not sure if this is right, but seems to give the relevant mnts */
 					if (!STREQLEN(mnt->mnt_fsname, "/dev", 4))
 						continue;
+					if (STREQLEN(mnt->mnt_fsname, "/dev/loop", 9))
+						continue;
 
 					len = strlen(mnt->mnt_dir);
 					if (len && mnt->mnt_dir[len - 1] != '/') {
diff --git a/source/blender/editors/undo/CMakeLists.txt b/source/blender/editors/undo/CMakeLists.txt
index 89832604ed8..2e3e73f34f4 100644
--- a/source/blender/editors/undo/CMakeLists.txt
+++ b/source/blender/editors/undo/CMakeLists.txt
@@ -22,6 +22,7 @@ set(INC
 	../include
 	../../blenkernel
 	../../blenlib
+	../../blenloader
 	../../blentranslation
 	../../makesdna
 	../../makesrna
diff --git a/source/blender/editors/undo/ed_undo.c b/source/blender/editors/undo/ed_undo.c
index e0a1faf04b8..c69fabdbd70 100644
--- a/source/blender/editors/undo/ed_undo.c
+++ b/source/blender/editors/undo/ed_undo.c
@@ -47,9 +47,12 @@
 #include "BKE_context.h"
 #include "BKE_global.h"
 #include "BKE_main.h"
+#include "BKE_report.h"
 #include "BKE_screen.h"
 #include "BKE_undo_system.h"
 
+#include "BLO_runtime.h"
+
 #include "ED_gpencil.h"
 #include "ED_render.h"
 #include "ED_screen.h"
@@ -101,7 +104,7 @@ void ED_undo_push(bContext *C, const char *str)
 }
 
 /* note: also check undo_history_exec() in bottom if you change notifiers */
-static int ed_undo_step(bContext *C, int step, const char *undoname)
+static int ed_undo_step(bContext *C, int step, const char *undoname, ReportList *reports)
 {
 	CLOG_INFO(&LOG, 1, "name='%s', step=%d", undoname, step);
 	wmWindowManager *wm = CTX_wm_manager(C);
@@ -111,6 +114,14 @@ static int ed_undo_step(bContext *C, int step, const char *undoname)
 	 * or they can just lead to freezing job in some other cases */
 	WM_jobs_kill_all(wm);
 
+	if (G.debug & G_DEBUG_IO) {
+		Main *bmain = CTX_data_main(C);
+		if (bmain->lock != NULL) {
+			BKE_report(reports, RPT_INFO, "Checking sanity of current .blend file *BEFORE* undo step.");
+			BLO_main_validate_libraries(bmain, reports);
+		}
+	}
+
 	/* TODO(campbell): undo_system: use undo system */
 	/* grease pencil can be can be used in plenty of spaces, so check it first */
 	if (ED_gpencil_session_active()) {
@@ -161,6 +172,14 @@ static int ed_undo_step(bContext *C, int step, const char *undoname)
 		wm->op_undo_depth--;
 	}
 
+	if (G.debug & G_DEBUG_IO) {
+		Main *bmain = CTX_data_main(C);
+		if (bmain->lock != NULL) {
+			BKE_report(reports, RPT_INFO, "Checking sanity of current .blend file *AFTER* undo step.");
+			BLO_main_validate_libraries(bmain, reports);
+		}
+	}
+
 	WM_event_add_notifier(C, NC_WINDOW, NULL);
 	WM_event_add_notifier(C, NC_WM | ND_UNDO, NULL);
 
@@ -182,11 +201,11 @@ void ED_undo_grouped_push(bContext *C, const char *str)
 
 void ED_undo_pop(bContext *C)
 {
-	ed_undo_step(C, 1, NULL);
+	ed_undo_step(C, 1, NULL, NULL);
 }
 void ED_undo_redo(bContext *C)
 {
-	ed_undo_step(C, -1, NULL);
+	ed_undo_step(C, -1, NULL, NULL);
 }
 
 void ED_undo_push_op(bContext *C, wmOperator *op)
@@ -208,7 +227,7 @@ void ED_undo_grouped_push_op(bContext *C, wmOperator *op)
 void ED_undo_pop_op(bContext *C, wmOperator *op)
 {
 	/* search back a couple of undo's, in case something else added pushes */
-	ed_undo_step(C, 0, op->type->name);
+	ed_undo_step(C, 0, op->type->name, op->reports);
 }
 
 /* name optionally, function used to check for operator redo panel */
@@ -237,11 +256,11 @@ UndoStack *ED_undo_stack_get(void)
 /** \name Undo, Undo Push & Redo Operators
  * \{ */
 
-static int ed_undo_exec(bContext *C, wmOperator *UNUSED(op))
+static int ed_undo_exec(bContext *C, wmOperator *op)
 {
 	/* "last operator" should disappear, later we can tie this with undo stack nicer */
 	WM_operator_stack_clear(CTX_wm_manager(C));
-	int ret = ed_undo_step(C, 1, NULL);
+	int ret = ed_undo_step(C, 1, NULL, op->reports);
 	if (ret & OPERATOR_FINISHED) {
 		/* Keep button under the cursor active. */
 		WM_event_add_mousemove(C);
@@ -257,9 +276,9 @@ static int ed_undo_push_exec(bContext *C, wmOperator *op)
 	return OPERATOR_FINISHED;
 }
 
-static int ed_redo_exec(bContext *C, wmOperator *UNUSED(op))
+static int ed_redo_exec(bContext *C, wmOperator *op)
 {
-	int ret = ed_undo_step(C, -1, NULL);
+	int ret = ed_undo_step(C, -1, NULL, op->reports);
 	if (ret & OPERATOR_FINISHED) {
 		/* Keep button under the cursor active. */
 		WM_event_add_mousemove(C);
diff --git a/source/blender/editors/uvedit/uvedit_parametrizer.c b/source/blender/editors/uvedit/uvedit_parametrizer.c
index 79e804725e5..352bfbf9a9e 100644
--- a/source/blender/editors/uvedit/uvedit_parametrizer.c
+++ b/source/blender/editors/uvedit/uvedit_parametrizer.c
@@ -34,6 +34,8 @@
 #include "BLI_heap.h"
 #include "BLI_boxpack_2d.h"
 #include "BLI_convexhull_2d.h"
+#include "BLI_polyfill_2d.h"
+#include "BLI_polyfill_2d_beautify.h"
 
 #include "uvedit_parametrizer.h"
 
@@ -219,6 +221,8 @@ enum PHandleState {
 typedef struct PHandle {
 	enum PHandleState state;
 	MemArena *arena;
+	MemArena *polyfill_arena;
+	Heap *polyfill_heap;
 
 	PChart *construction_chart;
 	PHash *hash_verts;
@@ -4119,6 +4123,8 @@ ParamHandle *param_construct_begin(void)
 	handle->construction_chart = p_chart_new(handle);
 	handle->state = PHANDLE_STATE_ALLOCATED;
 	handle->arena = BLI_memarena_new(MEM_SIZE_OPTIMAL(1 << 16), "param construct arena");
+	handle->polyfill_arena = BLI_memarena_new(BLI_MEMARENA_STD_BUFSIZE, "param polyfill arena");
+	handle->polyfill_heap = BLI_heap_new_ex(BLI_POLYFILL_ALLOC_NGON_RESERVE);
 	handle->aspx = 1.0f;
 	handle->aspy = 1.0f;
 	handle->do_aspect = false;
@@ -4162,82 +4168,71 @@ void param_delete(ParamHandle *handle)
 	}
 
 	BLI_memarena_free(phandle->arena);
+	BLI_memarena_free(phandle->polyfill_arena);
+	BLI_heap_free(phandle->polyfill_heap, NULL);
 	MEM_freeN(phandle);
 }
 
 static void p_add_ngon(ParamHandle *handle, ParamKey key, int nverts,
                        ParamKey *vkeys, float **co, float **uv,
-                       ParamBool *pin, ParamBool *select, const float normal[3])
+                       ParamBool *pin, ParamBool *select)
 {
-	int *boundary = BLI_array_alloca(boundary, nverts);
-
-	/* boundary vertex indexes */
-	for (int i = 0; i < nverts; i++) {
-		boundary[i] = i;
-	}
-
-	while (nverts > 2) {
-		float minangle = FLT_MAX;
-		float minshape = FLT_MAX;
-		int i, mini = 0;
-
-		/* find corner with smallest angle */
-		for (i = 0; i < nverts; i++) {
-			int v0 = boundary[(i + nverts - 1) % nverts];
-			int v1 = boundary[i];
-			int v2 = boundary[(i + 1) % nverts];
-			float angle = p_vec_angle(co[v0], co[v1], co[v2]);
-			float n[3];
-
-			normal_tri_v3(n, co[v0], co[v1], co[v2]);
+	/* Allocate memory for polyfill. */
+	PHandle *phandle = (PHandle *)handle;
+	MemArena *arena = phandle->polyfill_arena;
+	Heap *heap = phandle->polyfill_heap;
+	unsigned int nfilltri = nverts - 2;
+	unsigned int (*tris)[3] = BLI_memarena_alloc(arena, sizeof(*tris) * (size_t)nfilltri);
+	float (*projverts)[2] = BLI_memarena_alloc(arena, sizeof(*projverts) * (size_t)nverts);
 
-			if (normal && (dot_v3v3(n, normal) < 0.0f))
-				angle = (float)(2.0 * M_PI) - angle;
+	/* Calc normal, flipped: to get a positive 2d cross product. */
+	float normal[3];
+	zero_v3(normal);
 
-			float other_angle = p_vec_angle(co[v2], co[v0], co[v1]);
-			float shape = fabsf((float)M_PI - angle - 2.0f * other_angle);
+	const float *co_curr, *co_prev = co[nverts-1];
+	for (int j = 0; j < nverts; j++) {
+		co_curr = co[j];
+		add_newell_cross_v3_v3v3(normal, co_prev, co_curr);
+		co_prev = co_curr;
+	}
+	if (UNLIKELY(normalize_v3(normal) == 0.0f)) {
+		normal[2] = 1.0f;
+	}
 
-			if (fabsf(angle - minangle) < 0.01f) {
-				/* for nearly equal angles, try to get well shaped triangles */
-				if (shape < minshape) {
-					minangle = angle;
-					minshape = shape;
-					mini = i;
-				}
-			}
-			else if (angle < minangle) {
-				minangle = angle;
-				minshape = shape;
-				mini = i;
-			}
-		}
+	/* Project verts to 2d. */
+	float axis_mat[3][3];
+	axis_dominant_v3_to_m3_negate(axis_mat, normal);
+	for (int j = 0; j < nverts; j++) {
+		mul_v2_m3v3(projverts[j], axis_mat, co[j]);
+	}
 
-		/* add triangle in corner */
-		{
-			int v0 = boundary[(mini + nverts - 1) % nverts];
-			int v1 = boundary[mini];
-			int v2 = boundary[(mini + 1) % nverts];
+	BLI_polyfill_calc_arena(projverts, nverts, 1, tris, arena);
 
-			ParamKey tri_vkeys[3] = {vkeys[v0], vkeys[v1], vkeys[v2]};
-			float *tri_co[3] = {co[v0], co[v1], co[v2]};
-			float *tri_uv[3] = {uv[v0], uv[v1], uv[v2]};
-			ParamBool tri_pin[3] = {pin[v0], pin[v1], pin[v2]};
-			ParamBool tri_select[3] = {select[v0], select[v1], select[v2]};
+	/* Beautify helps avoid thin triangles that give numerical problems. */
+	BLI_polyfill_beautify(projverts, nverts, tris, arena, heap);
 
-			param_face_add(handle, key, 3, tri_vkeys, tri_co, tri_uv, tri_pin, tri_select, NULL);
-		}
+	/* Add triangles. */
+	for (int j = 0; j < nfilltri; j++) {
+		unsigned int *tri = tris[j];
+		unsigned int v0 = tri[0];
+		unsigned int v1 = tri[1];
+		unsigned int v2 = tri[2];
 
-		/* remove corner */
-		if (mini + 1 < nverts)
-			memmove(boundary + mini, boundary + mini + 1, (nverts - mini - 1) * sizeof(int));
+		ParamKey tri_vkeys[3] = {vkeys[v0], vkeys[v1], vkeys[v2]};
+		float *tri_co[3] = {co[v0], co[v1], co[v2]};
+		float *tri_uv[3] = {uv[v0], uv[v1], uv[v2]};
+		ParamBool tri_pin[3] = {pin[v0], pin[v1], pin[v2]};
+		ParamBool tri_select[3] = {select[v0], select[v1], select[v2]};
 
-		nverts--;
+		param_face_add(handle, key, 3, tri_vkeys, tri_co, tri_uv, tri_pin, tri_select);
 	}
+
+	BLI_memarena_clear(arena);
 }
 
 void param_face_add(ParamHandle *handle, ParamKey key, int nverts,
                     ParamKey *vkeys, float *co[4], float *uv[4],
-                    ParamBool *pin, ParamBool *select, float normal[3])
+                    ParamBool *pin, ParamBool *select)
 {
 	PHandle *phandle = (PHandle *)handle;
 
@@ -4247,7 +4242,7 @@ void param_face_add(ParamHandle *handle, ParamKey key, int nverts,
 
 	if (nverts > 4) {
 		/* ngon */
-		p_add_ngon(handle, key, nverts, vkeys, co, uv, pin, select, normal);
+		p_add_ngon(handle, key, nverts, vkeys, co, uv, pin, select);
 	}
 	else if (nverts == 4) {
 		/* quad */
diff --git a/source/blender/editors/uvedit/uvedit_parametrizer.h b/source/blender/editors/uvedit/uvedit_parametrizer.h
index 2714bc33769..e42944f3da4 100644
--- a/source/blender/editors/uvedit/uvedit_parametrizer.h
+++ b/source/blender/editors/uvedit/uvedit_parametrizer.h
@@ -62,8 +62,7 @@ void param_face_add(ParamHandle *handle,
                     float *co[4],
                     float *uv[4],
                     ParamBool *pin,
-                    ParamBool *select,
-                    float face_normal[3]);
+                    ParamBool *select);
 
 void param_edge_set_seam(ParamHandle *handle,
                          ParamKey *vkeys);
diff --git a/source/blender/editors/uvedit/uvedit_unwrap_ops.c b/source/blender/editors/uvedit/uvedit_unwrap_ops.c
index 582ba989997..088ed42b0c0 100644
--- a/source/blender/editors/uvedit/uvedit_unwrap_ops.c
+++ b/source/blender/editors/uvedit/uvedit_unwrap_ops.c
@@ -254,7 +254,7 @@ static void construct_param_handle_face_add(ParamHandle *handle, Scene *scene,
 		select[i] = uvedit_uv_select_test(scene, l, cd_loop_uv_offset);
 	}
 
-	param_face_add(handle, key, i, vkeys, co, uv, pin, select, efa->no);
+	param_face_add(handle, key, i, vkeys, co, uv, pin, select);
 }
 
 static ParamHandle *construct_param_handle(Scene *scene, Object *ob, BMesh *bm,
@@ -478,7 +478,7 @@ static ParamHandle *construct_param_handle_subsurfed(Scene *scene, Object *ob, B
 		texface_from_original_index(origFace, origVertIndices[mloop[2].v], &uv[2], &pin[2], &select[2], scene, cd_loop_uv_offset);
 		texface_from_original_index(origFace, origVertIndices[mloop[3].v], &uv[3], &pin[3], &select[3], scene, cd_loop_uv_offset);
 
-		param_face_add(handle, key, 4, vkeys, co, uv, pin, select, NULL);
+		param_face_add(handle, key, 4, vkeys, co, uv, pin, select);
 	}
 
 	/* these are calculated from original mesh too */
diff --git a/source/blender/gpu/CMakeLists.txt b/source/blender/gpu/CMakeLists.txt
index 0fde0e2a4a9..31cf2033714 100644
--- a/source/blender/gpu/CMakeLists.txt
+++ b/source/blender/gpu/CMakeLists.txt
@@ -62,28 +62,6 @@ set(SRC
 	intern/gpu_shader.c
 	intern/gpu_texture.c
 
-	shaders/gpu_shader_fx_lib.glsl
-	shaders/gpu_shader_fx_ssao_frag.glsl
-	shaders/gpu_shader_fx_dof_frag.glsl
-	shaders/gpu_shader_fx_dof_vert.glsl
-	shaders/gpu_shader_fx_dof_hq_frag.glsl
-	shaders/gpu_shader_fx_dof_hq_vert.glsl
-	shaders/gpu_shader_fx_dof_hq_geo.glsl
-	shaders/gpu_shader_fx_vert.glsl
-	shaders/gpu_shader_material.glsl
-	shaders/gpu_shader_sep_gaussian_blur_frag.glsl
-	shaders/gpu_shader_sep_gaussian_blur_vert.glsl
-	shaders/gpu_shader_basic_frag.glsl
-	shaders/gpu_shader_basic_vert.glsl
-	shaders/gpu_shader_basic_geom.glsl
-	shaders/gpu_shader_vertex.glsl
-	shaders/gpu_shader_vsm_store_frag.glsl
-	shaders/gpu_shader_vsm_store_vert.glsl
-	shaders/gpu_shader_fx_depth_resolve.glsl
-	shaders/gpu_shader_fire_frag.glsl
-	shaders/gpu_shader_smoke_frag.glsl
-	shaders/gpu_shader_smoke_vert.glsl
-
 	GPU_basic_shader.h
 	GPU_buffers.h
 	GPU_compositing.h
diff --git a/source/blender/makesdna/DNA_modifier_types.h b/source/blender/makesdna/DNA_modifier_types.h
index ad5f3768b93..2adf4a02d22 100644
--- a/source/blender/makesdna/DNA_modifier_types.h
+++ b/source/blender/makesdna/DNA_modifier_types.h
@@ -192,8 +192,9 @@ typedef struct MaskModifierData {
 	struct Object *ob_arm;  /* armature to use to in place of hardcoded vgroup */
 	char vgroup[64];        /* name of vertex group to use to mask, MAX_VGROUP_NAME */
 
-	int mode;               /* using armature or hardcoded vgroup */
-	int flag;               /* flags for various things */
+	short mode;               /* using armature or hardcoded vgroup */
+	short flag;               /* flags for various things */
+	float threshold;
 } MaskModifierData;
 
 /* Mask Modifier -> mode */
diff --git a/source/blender/makesrna/intern/rna_modifier.c b/source/blender/makesrna/intern/rna_modifier.c
index 1a05b6e5e00..f9b848744d6 100644
--- a/source/blender/makesrna/intern/rna_modifier.c
+++ b/source/blender/makesrna/intern/rna_modifier.c
@@ -3267,6 +3267,13 @@ static void rna_def_modifier_mask(BlenderRNA *brna)
 	RNA_def_property_boolean_sdna(prop, NULL, "flag", MOD_MASK_INV);
 	RNA_def_property_ui_text(prop, "Invert", "Use vertices that are not part of region defined");
 	RNA_def_property_update(prop, 0, "rna_Modifier_update");
+
+	prop = RNA_def_property(srna, "threshold", PROP_FLOAT, PROP_FACTOR);
+	RNA_def_property_float_sdna(prop, NULL, "threshold");
+	RNA_def_property_range(prop, 0.0, 1.0);
+	RNA_def_property_ui_range(prop, 0, 1, 0.1, 3);
+	RNA_def_property_ui_text(prop, "Threshold", "Weights over this threshold remain");
+	RNA_def_property_update(prop, 0, "rna_Modifier_update");
 }
 
 static void rna_def_modifier_simpledeform(BlenderRNA *brna)
diff --git a/source/blender/modifiers/intern/MOD_mask.c b/source/blender/modifiers/intern/MOD_mask.c
index 93332fb0455..5246f8cd35a 100644
--- a/source/blender/modifiers/intern/MOD_mask.c
+++ b/source/blender/modifiers/intern/MOD_mask.c
@@ -196,7 +196,7 @@ static DerivedMesh *applyModifier(
 			for (j = 0; j < dv->totweight; j++, dw++) {
 				if (dw->def_nr < defbase_tot) {
 					if (bone_select_array[dw->def_nr]) {
-						if (dw->weight != 0.0f) {
+						if (dw->weight > mmd->threshold) {
 							found = true;
 							break;
 						}
@@ -228,7 +228,7 @@ static DerivedMesh *applyModifier(
 
 		/* add vertices which exist in vertexgroup into ghash for filtering */
 		for (i = 0, dv = dvert; i < maxVerts; i++, dv++) {
-			const bool found = defvert_find_weight(dv, defgrp_index) != 0.0f;
+			const bool found = defvert_find_weight(dv, defgrp_index) > mmd->threshold;
 			if (found_test != found) {
 				continue;
 			}
diff --git a/source/blender/python/intern/bpy_app.c b/source/blender/python/intern/bpy_app.c
index d519b5f8aef..928e14c09bf 100644
--- a/source/blender/python/intern/bpy_app.c
+++ b/source/blender/python/intern/bpy_app.c
@@ -364,6 +364,7 @@ static PyGetSetDef bpy_app_getsets[] = {
 	{(char *)"debug_depsgraph_pretty", bpy_app_debug_get, bpy_app_debug_set, (char *)bpy_app_debug_doc, (void *)G_DEBUG_DEPSGRAPH_PRETTY},
 	{(char *)"debug_simdata",   bpy_app_debug_get, bpy_app_debug_set, (char *)bpy_app_debug_doc, (void *)G_DEBUG_SIMDATA},
 	{(char *)"debug_gpumem",    bpy_app_debug_get, bpy_app_debug_set, (char *)bpy_app_debug_doc, (void *)G_DEBUG_GPU_MEM},
+	{(char *)"debug_io",        bpy_app_debug_get, bpy_app_debug_set, (char *)bpy_app_debug_doc, (void *)G_DEBUG_IO},
 
 	{(char *)"binary_path_python", bpy_app_binary_path_python_get, NULL, (char *)bpy_app_binary_path_python_doc, NULL},
 
diff --git a/source/blender/windowmanager/WM_api.h b/source/blender/windowmanager/WM_api.h
index 2ff416c0ba6..5ee7763272b 100644
--- a/source/blender/windowmanager/WM_api.h
+++ b/source/blender/windowmanager/WM_api.h
@@ -55,6 +55,7 @@ struct wmGesture;
 struct wmJob;
 struct wmOperatorType;
 struct wmOperator;
+struct wmPaintCursor;
 struct rcti;
 struct PointerRNA;
 struct PropertyRNA;
@@ -133,15 +134,17 @@ void		WM_cursor_grab_enable(struct wmWindow *win, bool wrap, bool hide, int boun
 void		WM_cursor_grab_disable(struct wmWindow *win, const int mouse_ungrab_xy[2]);
 void		WM_cursor_time		(struct wmWindow *win, int nr);
 
-void *WM_paint_cursor_activate(
+struct wmPaintCursor *WM_paint_cursor_activate(
         struct wmWindowManager *wm,
         bool (*poll)(struct bContext *C),
         void (*draw)(struct bContext *C, int, int, void *customdata),
         void *customdata);
 
-void		WM_paint_cursor_end(struct wmWindowManager *wm, void *handle);
+bool		WM_paint_cursor_end(struct wmWindowManager *wm, struct wmPaintCursor *handle);
+void       *WM_paint_cursor_customdata_get(struct wmPaintCursor *pc);
 void		WM_paint_cursor_tag_redraw(struct wmWindow *win, struct ARegion *ar);
 
+
 void		WM_cursor_warp		(struct wmWindow *win, int x, int y);
 void		WM_cursor_compatible_xy(wmWindow *win, int *x, int *y);
 float		WM_cursor_pressure	(const struct wmWindow *win);
diff --git a/source/blender/windowmanager/intern/wm_files.c b/source/blender/windowmanager/intern/wm_files.c
index 796103677bf..87ca7310971 100644
--- a/source/blender/windowmanager/intern/wm_files.c
+++ b/source/blender/windowmanager/intern/wm_files.c
@@ -1128,12 +1128,12 @@ bool write_crash_blend(void)
 /**
  * \see #wm_homefile_write_exec wraps #BLO_write_file in a similar way.
  */
-static int wm_file_write(bContext *C, const char *filepath, int fileflags, ReportList *reports)
+static bool wm_file_write(bContext *C, const char *filepath, int fileflags, ReportList *reports)
 {
 	Main *bmain = CTX_data_main(C);
 	Library *li;
 	int len;
-	int ret = -1;
+	int ok = false;
 	BlendThumbnail *thumb, *main_thumb;
 	ImBuf *ibuf_thumb = NULL;
 
@@ -1141,18 +1141,18 @@ static int wm_file_write(bContext *C, const char *filepath, int fileflags, Repor
 
 	if (len == 0) {
 		BKE_report(reports, RPT_ERROR, "Path is empty, cannot save");
-		return ret;
+		return ok;
 	}
 
 	if (len >= FILE_MAX) {
 		BKE_report(reports, RPT_ERROR, "Path too long, cannot save");
-		return ret;
+		return ok;
 	}
 
 	/* Check if file write permission is ok */
 	if (BLI_exists(filepath) && !BLI_file_is_writable(filepath)) {
 		BKE_reportf(reports, RPT_ERROR, "Cannot save blend file, path '%s' is not writable", filepath);
-		return ret;
+		return ok;
 	}
 
 	/* note: used to replace the file extension (to ensure '.blend'),
@@ -1163,7 +1163,7 @@ static int wm_file_write(bContext *C, const char *filepath, int fileflags, Repor
 	for (li = bmain->library.first; li; li = li->id.next) {
 		if (BLI_path_cmp(li->filepath, filepath) == 0) {
 			BKE_reportf(reports, RPT_ERROR, "Cannot overwrite used library '%.240s'", filepath);
-			return ret;
+			return ok;
 		}
 	}
 
@@ -1226,7 +1226,8 @@ static int wm_file_write(bContext *C, const char *filepath, int fileflags, Repor
 			ibuf_thumb = IMB_thumb_create(filepath, THB_LARGE, THB_SOURCE_BLEND, ibuf_thumb);
 		}
 
-		ret = 0;  /* Success. */
+		/* Success. */
+		ok = true;
 	}
 
 	if (ibuf_thumb) {
@@ -1238,7 +1239,7 @@ static int wm_file_write(bContext *C, const char *filepath, int fileflags, Repor
 
 	WM_cursor_wait(0);
 
-	return ret;
+	return ok;
 }
 
 /************************ autosave ****************************/
@@ -1415,6 +1416,7 @@ void WM_file_tag_modified(void)
 
 /**
  * \see #wm_file_write wraps #BLO_write_file in a similar way.
+ * \return success.
  */
 static int wm_homefile_write_exec(bContext *C, wmOperator *op)
 {
@@ -2092,7 +2094,7 @@ static int wm_save_as_mainfile_exec(bContext *C, wmOperator *op)
 {
 	Main *bmain = CTX_data_main(C);
 	char path[FILE_MAX];
-	int fileflags;
+	const bool is_save_as = (op->type->invoke == wm_save_as_mainfile_invoke);
 
 	save_set_compress(op);
 
@@ -2104,7 +2106,8 @@ static int wm_save_as_mainfile_exec(bContext *C, wmOperator *op)
 		wm_filepath_default(path);
 	}
 
-	fileflags = G.fileflags & ~G_FILE_USERPREFS;
+	const int fileflags_orig = G.fileflags;
+	int fileflags = G.fileflags & ~G_FILE_USERPREFS;
 
 	/* set compression flag */
 	SET_FLAG_FROM_TEST(
@@ -2119,12 +2122,22 @@ static int wm_save_as_mainfile_exec(bContext *C, wmOperator *op)
 	         RNA_boolean_get(op->ptr, "copy")),
 	        G_FILE_SAVE_COPY);
 
-	if (wm_file_write(C, path, fileflags, op->reports) != 0)
+	const bool ok = wm_file_write(C, path, fileflags, op->reports);
+
+	if ((op->flag & OP_IS_INVOKE) == 0) {
+		/* OP_IS_INVOKE is set when the operator is called from the GUI.
+		 * If it is not set, the operator is called from a script and
+		 * shouldn't influence G.fileflags. */
+		G.fileflags = fileflags_orig;
+	}
+
+	if (ok == false) {
 		return OPERATOR_CANCELLED;
+	}
 
 	WM_event_add_notifier(C, NC_WM | ND_FILESAVE, NULL);
 
-	if (RNA_boolean_get(op->ptr, "exit")) {
+	if (!is_save_as && RNA_boolean_get(op->ptr, "exit")) {
 		wm_exit_schedule_delayed(C);
 	}
 
diff --git a/source/blender/windowmanager/intern/wm_init_exit.c b/source/blender/windowmanager/intern/wm_init_exit.c
index 185cf3fad4f..d254a83fcc2 100644
--- a/source/blender/windowmanager/intern/wm_init_exit.c
+++ b/source/blender/windowmanager/intern/wm_init_exit.c
@@ -188,7 +188,10 @@ void WM_init(bContext *C, int argc, const char **argv)
 	ED_node_init_butfuncs();
 
 	BLF_init(); /* Please update source/gamengine/GamePlayer/GPG_ghost.cpp if you change this */
+
 	BLT_lang_init();
+	/* Must call first before doing any .blend file reading, since versionning code may create new IDs... See T57066. */
+	BLT_lang_set(NULL);
 
 	/* reports cant be initialized before the wm,
 	 * but keep before file reading, since that may report errors */
@@ -197,6 +200,7 @@ void WM_init(bContext *C, int argc, const char **argv)
 	/* get the default database, plus a wm */
 	wm_homefile_read(C, NULL, G.factory_startup, false, true, NULL, WM_init_state_app_template_get());
 
+	/* Call again to set from userpreferences... */
 	BLT_lang_set(NULL);
 
 	if (!G.background) {
diff --git a/source/blender/windowmanager/intern/wm_operators.c b/source/blender/windowmanager/intern/wm_operators.c
index 45084980727..2008d388ad3 100644
--- a/source/blender/windowmanager/intern/wm_operators.c
+++ b/source/blender/windowmanager/intern/wm_operators.c
@@ -1813,7 +1813,7 @@ static void WM_OT_console_toggle(wmOperatorType *ot)
  * - draw(bContext): drawing callback for paint cursor
  */
 
-void *WM_paint_cursor_activate(
+wmPaintCursor *WM_paint_cursor_activate(
         wmWindowManager *wm, bool (*poll)(bContext *C),
         wmPaintCursorDraw draw, void *customdata)
 {
@@ -1828,7 +1828,7 @@ void *WM_paint_cursor_activate(
 	return pc;
 }
 
-void WM_paint_cursor_end(wmWindowManager *wm, void *handle)
+bool WM_paint_cursor_end(wmWindowManager *wm, wmPaintCursor *handle)
 {
 	wmPaintCursor *pc;
 
@@ -1836,9 +1836,15 @@ void WM_paint_cursor_end(wmWindowManager *wm, void *handle)
 		if (pc == (wmPaintCursor *)handle) {
 			BLI_remlink(&wm->paintcursors, pc);
 			MEM_freeN(pc);
-			return;
+			return true;
 		}
 	}
+	return false;
+}
+
+void *WM_paint_cursor_customdata_get(wmPaintCursor *pc)
+{
+	return pc->customdata;
 }
 
 /* *********************** radial control ****************** */
diff --git a/source/creator/CMakeLists.txt b/source/creator/CMakeLists.txt
index e6772a003f5..fb02a7c871c 100644
--- a/source/creator/CMakeLists.txt
+++ b/source/creator/CMakeLists.txt
@@ -753,6 +753,19 @@ elseif(WIN32)
 				DESTINATION ${BLENDER_VERSION}/python/bin
 				CONFIGURATIONS Debug
 			)
+			if(WINDOWS_PYTHON_DEBUG)
+				install(
+					FILES ${LIBDIR}/python/lib/python${_PYTHON_VERSION_NO_DOTS}.pdb
+					DESTINATION "."
+					CONFIGURATIONS Release;RelWithDebInfo;MinSizeRel
+				)
+
+				install(
+					FILES ${LIBDIR}/python/lib/python${_PYTHON_VERSION_NO_DOTS}_d.pdb
+					DESTINATION "."
+					CONFIGURATIONS Debug
+				)
+			endif()
 		endif()
 
 		unset(_PYTHON_VERSION_NO_DOTS)
@@ -1031,4 +1044,9 @@ if(WIN32 AND NOT WITH_PYTHON_MODULE)
 		COMPONENT Blender
 		DESTINATION "."
 	)
+	set_target_properties(
+		blender
+		PROPERTIES
+			VS_USER_PROPS "blender.Cpp.user.props"
+	)
 endif()