Cycles: Added Cryptomatte output.

This allows for extra output passes that encode automatic object and material masks
for the entire scene. It is an implementation of the Cryptomatte standard as
introduced by Psyop. A good future extension would be to add a manifest to the
export and to do plenty of testing to ensure that it is fully compatible with other
renderers and compositing programs that use Cryptomatte.

Internally, it adds the ability for Cycles to have several passes of the same type
that are distinguished by their name.

Differential Revision: https://developer.blender.org/D3538

32 files changed, 810 insertions, 46 deletions

diff --git a/intern/cycles/blender/addon/engine.py b/intern/cycles/blender/addon/engine.py
index 16ec7bc314b..2cdeb97a32d 100644
--- a/intern/cycles/blender/addon/engine.py
+++ b/intern/cycles/blender/addon/engine.py
@@ -254,6 +254,17 @@ 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_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:
         engine.register_pass(scene, srl, "Noisy Image", 3, "RGBA", 'COLOR')
         if crl.denoising_store_passes:
diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index 80b83c94012..848b76eb02f 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -1339,7 +1339,36 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
             default=False,
             update=update_render_passes,
         )
-
+        cls.use_pass_crypto_object = BoolProperty(
+                name="Cryptomatte Object",
+                description="Cryptomatte Object pass",
+                default=False,
+                update=update_render_passes,
+                )
+        cls.use_pass_crypto_material = BoolProperty(
+                name="Cryptomatte Material",
+                description="Cryptomatte Material pass",
+                default=False,
+                update=update_render_passes,
+                )
+        cls.use_pass_crypto_asset = BoolProperty(
+                name="Cryptomatte Asset",
+                description="Cryptomatte Asset pass",
+                default=False,
+                update=update_render_passes,
+                )
+        cls.pass_crypto_depth = IntProperty(
+                name="Cryptomatte Levels",
+                description="Describes how many unique IDs per pixel are written to Cryptomatte",
+                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 5edbcb19672..6f11d3c313d 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -563,6 +563,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"
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..e9e14a9b6c9 100644
--- a/intern/cycles/blender/blender_session.cpp
+++ b/intern/cycles/blender/blender_session.cpp
@@ -405,7 +405,7 @@ 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");
@@ -700,7 +700,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 +719,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 8ae52beb1c1..076734d105f 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;
@@ -549,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) {
@@ -636,6 +641,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 7eb73dea3ef..eb816e1fdd0 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;
@@ -691,6 +699,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);
 				}
@@ -700,6 +711,9 @@ public:
 
 			task.update_progress(&tile, tile.w*tile.h);
 		}
+		if(use_coverage) {
+			coverage.finalize();
+		}
 	}
 
 	void denoise(DenoisingTask& denoising, RenderTile &tile)
@@ -760,7 +774,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 b48ed649a8c..08efede36df 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -96,6 +96,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 cfe17e63627..0eb8ce2cf8b 100644
--- a/intern/cycles/kernel/geom/geom_object.h
+++ b/intern/cycles/kernel/geom/geom_object.h
@@ -304,6 +304,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_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 e93100a6442..f46b06f87f9 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -53,6 +53,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
 
@@ -415,6 +416,7 @@ typedef enum PassType {
 	PASS_RAY_BOUNCES,
 #endif
 	PASS_RENDER_TIME,
+	PASS_CRYPTOMATTE,
 	PASS_CATEGORY_MAIN_END = 31,
 
 	PASS_MIST = 32,
@@ -443,6 +445,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,
@@ -1260,17 +1270,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. */
@@ -1460,7 +1473,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);
 
@@ -1540,7 +1557,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/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 e6021f4b37d..dd20efb3dde 100644
--- a/intern/cycles/render/buffers.cpp
+++ b/intern/cycles/render/buffers.cpp
@@ -233,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;
@@ -249,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();
 
@@ -385,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/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..c1f81e61b72
--- /dev/null
+++ b/intern/cycles/util/util_murmurhash.cpp
@@ -0,0 +1,125 @@
+/*
+ * 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 "util/util_algorithm.h"
+#include "util/util_murmurhash.h"
+
+#if defined(_MSC_VER)
+#  include <stdlib.h>
+#  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__ */