Cycles: Added Cryptomatte output.
authorStefan Werner <stefan.werner@tangent-animation.com>
Sun, 28 Oct 2018 09:37:41 +0000 (05:37 -0400)
committerStefan Werner <stefan.werner@tangent-animation.com>
Sun, 28 Oct 2018 09:37:41 +0000 (05:37 -0400)
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:
intern/cycles/blender/addon/engine.py
intern/cycles/blender/addon/properties.py
intern/cycles/blender/addon/ui.py
intern/cycles/blender/blender_object.cpp
intern/cycles/blender/blender_session.cpp
intern/cycles/blender/blender_sync.cpp
intern/cycles/blender/blender_sync.h
intern/cycles/device/device_cpu.cpp
intern/cycles/kernel/CMakeLists.txt
intern/cycles/kernel/geom/geom_object.h
intern/cycles/kernel/kernel_globals.h
intern/cycles/kernel/kernel_id_passes.h [new file with mode: 0644]
intern/cycles/kernel/kernel_passes.h
intern/cycles/kernel/kernel_shader.h
intern/cycles/kernel/kernel_types.h
intern/cycles/kernel/kernels/cuda/kernel.cu
intern/cycles/kernel/kernels/opencl/kernel.cl
intern/cycles/kernel/split/kernel_buffer_update.h
intern/cycles/render/CMakeLists.txt
intern/cycles/render/buffers.cpp
intern/cycles/render/buffers.h
intern/cycles/render/coverage.cpp [new file with mode: 0644]
intern/cycles/render/coverage.h [new file with mode: 0644]
intern/cycles/render/film.cpp
intern/cycles/render/film.h
intern/cycles/render/object.cpp
intern/cycles/render/object.h
intern/cycles/render/shader.cpp
intern/cycles/util/CMakeLists.txt
intern/cycles/util/util_atomic.h
intern/cycles/util/util_murmurhash.cpp [new file with mode: 0644]
intern/cycles/util/util_murmurhash.h [new file with mode: 0644]

index 16ec7bc314bf48847a0c66f4a3550209d4048111..2cdeb97a32d8e9dd49dcc87b08a70f6139e49b9e 100644 (file)
@@ -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:
index 80b83c940123f690e2416d64716198e899ad620f..848b76eb02f504a4589b0fbf74c7174266586142 100644 (file)
@@ -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
index 5edbcb1967276f44f5e6a99c16caf9e478ca8eba..6f11d3c313d3e20796046b01bdeee30e7fa2244b 100644 (file)
@@ -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"
index 0fab9ab35316e78bb31e4fa023d699c8dc9a0c48..a05c982b367f5a512add5b12d5c1fdf97cc8584c 100644 (file)
@@ -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 */
index a07131d04aebdf8284b1819c706450547b43dcf9..e9e14a9b6c944b9432e968d4b2c80397b77c5750 100644 (file)
@@ -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]);
        }
 
index 8ae52beb1c127a1c61b627bfbc068ad160b63097..076734d105f481262dab9deefa95b3390506770c 100644 (file)
@@ -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;
 }
 
index 5e63f76033d29760ea2178b97bbdab9527ece8ff..eb84bedb1184e4c2e489e4791d03aa742de15b8e 100644 (file)
@@ -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,
index 7eb73dea3ef610300d447a4078a6a5ca1ef0333c..eb816e1fdd0805158722e1c7c5246c562c37136d 100644 (file)
@@ -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);
                        }
 
index b48ed649a8c48bdc79a3599d0a3d05fad4b9050c..08efede36df7a8d2c3c425c91466ca2f15d2f155 100644 (file)
@@ -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
index cfe17e6362767d62c745822951af534069eed587..0eb8ce2cf8b2a233e4cfe448a03e7016f036e25d 100644 (file)
@@ -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)
index 74cfacb5bc1d6fbc37d2ecac1449a3c7757dbdf9..37402f428634d5b68c817ee1a6f13a4fe769da03 100644 (file)
@@ -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 (file)
index 0000000..486c61d
--- /dev/null
@@ -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
index 458aa6c2a972c292743fc4a873200280ae5e4c51..e256a1819ed6f3e3efbacb4f8c4507829a3c7403 100644 (file)
  * 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))
index e834b701f963ff3d7fe6b5a36b2017ecff04b7cb..af883aa715bc98f87bb2247308b607bf6babe120 100644 (file)
@@ -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
index e93100a64428e3ece8f15c7c58a58154937d722b..f46b06f87f93b987009560c50c784767fbd798c1 100644 (file)
@@ -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;
index 8a180a509e8f93173e107298de6b5f85cd3d2657..af311027f785f1e380433a3314efd5bf93929b86 100644 (file)
@@ -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
 
index 63128d0aecf37871ab8204f92ea75b2ce7f0d380..de1f5088629957d653ceead5c2695d562f19b451 100644 (file)
@@ -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__ */
index 180c0b57077dc984b89e534c87463ca5083e52a7..18eec6372f1699d50f7fde2a962e098ceca06a6c 100644 (file)
@@ -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;
index 7d2220f37f9b334db8e51eb9575ce71f435725b9..c0ce7368771a79b1ad0922b7c9970b2d40c4ca10 100644 (file)
@@ -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
index e6021f4b37d628fa3b87b3ac74ed7e03f8c37d51..dd20efb3ddef8caca0677f3a48489dd4e4567bff 100644 (file)
@@ -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]);
index 1b06ffe33a6a882552fc0bb4f16630323c581bf8..a8f019dddd6537b5fe814af9abd63ac7584f50f0 100644 (file)
@@ -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 (file)
index 0000000..72ef4cd
--- /dev/null
@@ -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 (file)
index 0000000..16176ce
--- /dev/null
@@ -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__ */
index 8f3596ade58bcefe0e159432453ab78ef01c267f..d0f15496e50a40d9b0d59ef46d3537cde49e6e3e 100644 (file)
@@ -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);
index 6ab2eea79b8212ee8b9f103239cffa9c8d296570..57f1bf4eb64c83abaa1c88ded8b26b74401fd03f 100644 (file)
@@ -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);
 };
 
index e3f35c366d62109708f893f82794f8e3fd3b46c4..a56a8a6ec58e8a85691dd5837c068114cf922380 100644 (file)
@@ -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) {
index b80c4aef70ba80031b7031070ae1e615f0fed74f..bd44b35aba312600dcb1af1d0dc8c7593b070cd8 100644 (file)
@@ -48,6 +48,7 @@ public:
        BoundBox bounds;
        uint random_id;
        int pass_id;
+       ustring asset_name;
        vector<ParamValue> attributes;
        uint visibility;
        array<Transform> motion;
index ac605305b9467b76160e57ff774bba9f5331211e..8d0cec7b14e83c2d20f184ec7239bb0cd5850fd4 100644 (file)
@@ -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;
index 291f9a9fcaefe8fcffdcac1e00f49f95344a0853..4f623c5dfb7cf3a0caf22f7c50318632e9cd0db2 100644 (file)
@@ -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
index f3c7ae546a0ace4aec938eb394da306255b86fdd..e17e99d0acd981297d9558d6a7a32ba26e3a373e 100644 (file)
@@ -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 (file)
index 0000000..c1f81e6
--- /dev/null
@@ -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 (file)
index 0000000..824ed59
--- /dev/null
@@ -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__ */