sandey/blender
1
0
Fork 0
forked from bartvdbraak/blender
Code Pull Requests Activity

Merge branch 'master' into blender2.8

Browse Source
This commit is contained in:
Brecht Van Lommel 2018-10-28 16:41:30 +01:00
commit 046735d751
41 changed files with 1037 additions and 125 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
build_files/cmake/macros.cmake
View File

@ -672,10 +672,10 @@ function(SETUP_BLENDER_SORTED_LIBS)
bf_intern_mikktspace
bf_intern_dualcon
bf_intern_cycles
cycles_device
cycles_render
cycles_graph
cycles_bvh
cycles_device
cycles_kernel
cycles_util
cycles_subd

17
intern/cycles/blender/addon/engine.py
View File

@ -257,7 +257,20 @@ def register_passes(engine, scene, srl):
if crl.use_pass_volume_indirect: engine.register_pass(scene, srl, "VolumeInd", 3, "RGB", 'COLOR')
cscene = scene.cycles
if crl.use_denoising and crl.denoising_store_passes and not cscene.use_progressive_refine:
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:
engine.register_pass(scene, srl, "Denoising Normal", 3, "XYZ", 'VECTOR')
engine.register_pass(scene, srl, "Denoising Normal Variance", 3, "XYZ", 'VECTOR')
engine.register_pass(scene, srl, "Denoising Albedo", 3, "RGB", 'COLOR')
@ -266,9 +279,7 @@ def register_passes(engine, scene, srl):
engine.register_pass(scene, srl, "Denoising Depth Variance", 1, "Z", 'VALUE')
engine.register_pass(scene, srl, "Denoising Shadow A", 3, "XYV", 'VECTOR')
engine.register_pass(scene, srl, "Denoising Shadow B", 3, "XYV", 'VECTOR')
engine.register_pass(scene, srl, "Denoising Image", 3, "RGB", 'COLOR')
engine.register_pass(scene, srl, "Denoising Image Variance", 3, "RGB", 'COLOR')
clean_options = ("denoising_diffuse_direct", "denoising_diffuse_indirect",
"denoising_glossy_direct", "denoising_glossy_indirect",
"denoising_transmission_direct", "denoising_transmission_indirect",

30
intern/cycles/blender/addon/properties.py
View File

@ -1347,6 +1347,36 @@ class CyclesRenderLayerSettings(bpy.types.PropertyGroup):
default=False,
update=update_render_passes,
)
use_pass_crypto_object: BoolProperty(
name="Cryptomatte Object",
description="Cryptomatte Object pass",
default=False,
update=update_render_passes,
)
use_pass_crypto_material: BoolProperty(
name="Cryptomatte Material",
description="Cryptomatte Material pass",
default=False,
update=update_render_passes,
)
use_pass_crypto_asset: BoolProperty(
name="Cryptomatte Asset",
description="Cryptomatte Asset pass",
default=False,
update=update_render_passes,
)
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,
)
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 register(cls):

11
intern/cycles/blender/addon/ui.py
View File

@ -817,6 +817,17 @@ class CYCLES_RENDER_PT_layer_passes(CyclesButtonsPanel, Panel):
col.prop(cycles_view_layer, "pass_debug_bvh_intersections")
col.prop(cycles_view_layer, "pass_debug_ray_bounces")
layout.label("Cryptomatte:")
row = layout.row(align=True)
row.prop(cycles_view_layer, "use_pass_crypto_object", text="Object", toggle=True)
row.prop(cycles_view_layer, "use_pass_crypto_material", text="Material", toggle=True)
row.prop(cycles_view_layer, "use_pass_crypto_asset", text="Asset", toggle=True)
row = layout.row(align=True)
row.prop(cycles_view_layer, "pass_crypto_depth")
row = layout.row(align=True)
row.active = use_cpu(context)
row.prop(cycles_view_layer, "pass_crypto_accurate", text="Accurate Mode")
class CYCLES_RENDER_PT_denoising(CyclesButtonsPanel, Panel):
bl_label = "Denoising"

17
intern/cycles/blender/blender_object.cpp
View File

@ -424,6 +424,23 @@ Object *BlenderSync::sync_object(BL::Depsgraph& b_depsgraph,
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 */

6
intern/cycles/blender/blender_session.cpp
View File

@ -407,7 +407,7 @@ void BlenderSession::render(BL::Depsgraph& b_depsgraph_)
BL::RenderLayer b_rlay = *b_single_rlay;
/* add passes */
array<Pass> passes = sync->sync_render_passes(b_rlay, b_view_layer, session_params);
vector<Pass> passes = sync->sync_render_passes(b_rlay, b_view_layer, session_params);
buffer_params.passes = passes;
PointerRNA crl = RNA_pointer_get(&b_view_layer.ptr, "cycles");
@ -711,7 +711,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);
@ -730,7 +730,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]);
}

74
intern/cycles/blender/blender_sync.cpp
View File

@ -40,6 +40,8 @@
CCL_NAMESPACE_BEGIN
static const char *cryptomatte_prefix = "Crypto";
/* Constructor */
BlenderSync::BlenderSync(BL::RenderEngine& b_engine,
@ -449,6 +451,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;
@ -457,6 +462,9 @@ PassType BlenderSync::get_pass_type(BL::RenderPass& b_pass)
int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
{
string name = b_pass.name();
if(name == "Noisy Image") return DENOISING_PASS_COLOR;
if(name.substr(0, 10) != "Denoising ") {
return -1;
}
@ -471,7 +479,6 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
MAP_PASS("Depth Variance", DENOISING_PASS_DEPTH_VAR);
MAP_PASS("Shadow A", DENOISING_PASS_SHADOW_A);
MAP_PASS("Shadow B", DENOISING_PASS_SHADOW_B);
MAP_PASS("Image", DENOISING_PASS_COLOR);
MAP_PASS("Image Variance", DENOISING_PASS_COLOR_VAR);
MAP_PASS("Clean", DENOISING_PASS_CLEAN);
#undef MAP_PASS
@ -479,11 +486,11 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
return -1;
}
array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
BL::ViewLayer& b_view_layer,
const SessionParams &session_params)
{
array<Pass> passes;
vector<Pass> passes;
Pass::add(PASS_COMBINED, passes);
if(!session_params.device.advanced_shading) {
@ -505,20 +512,7 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
scene->film->denoising_flags = 0;
PointerRNA crp = RNA_pointer_get(&b_view_layer.ptr, "cycles");
if(get_boolean(crp, "denoising_store_passes") &&
get_boolean(crp, "use_denoising"))
{
b_engine.add_pass("Denoising Normal", 3, "XYZ", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Albedo", 3, "RGB", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Albedo Variance", 3, "RGB", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Depth", 1, "Z", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Depth Variance", 1, "Z", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Shadow A", 3, "XYV", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Shadow B", 3, "XYV", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Image", 3, "RGB", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Image Variance", 3, "RGB", b_view_layer.name().c_str());
if(get_boolean(crp, "use_denoising")) {
#define MAP_OPTION(name, flag) if(!get_boolean(crp, name)) scene->film->denoising_flags |= flag;
MAP_OPTION("denoising_diffuse_direct", DENOISING_CLEAN_DIFFUSE_DIR);
MAP_OPTION("denoising_diffuse_indirect", DENOISING_CLEAN_DIFFUSE_IND);
@ -530,10 +524,23 @@ array<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
MAP_OPTION("denoising_subsurface_indirect", DENOISING_CLEAN_SUBSURFACE_IND);
#undef MAP_OPTION
b_engine.add_pass("Noisy Image", 4, "RGBA", b_view_layer.name().c_str());
if(get_boolean(crp, "denoising_store_passes")) {
b_engine.add_pass("Denoising Normal", 3, "XYZ", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Albedo", 3, "RGB", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Albedo Variance", 3, "RGB", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Depth", 1, "Z", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Depth Variance", 1, "Z", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Shadow A", 3, "XYV", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Shadow B", 3, "XYV", b_view_layer.name().c_str());
b_engine.add_pass("Denoising Image Variance", 3, "RGB", b_view_layer.name().c_str());
if(scene->film->denoising_flags & DENOISING_CLEAN_ALL_PASSES) {
b_engine.add_pass("Denoising Clean", 3, "RGB", b_view_layer.name().c_str());
}
}
}
#ifdef __KERNEL_DEBUG__
if(get_boolean(crp, "pass_debug_bvh_traversed_nodes")) {
b_engine.add_pass("Debug BVH Traversed Nodes", 1, "X", b_view_layer.name().c_str());
@ -565,6 +572,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_view_layer.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_view_layer.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_view_layer.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;
}

2
intern/cycles/blender/blender_sync.h
View File

@ -67,7 +67,7 @@ public:
int width, int height,
void **python_thread_state);
void sync_view_layer(BL::SpaceView3D& b_v3d, BL::ViewLayer& b_view_layer);
array<Pass> sync_render_passes(BL::RenderLayer& b_render_layer,
vector<Pass> sync_render_passes(BL::RenderLayer& b_render_layer,
BL::ViewLayer& b_view_layer,
const SessionParams &session_params);
void sync_integrator();

15
intern/cycles/device/device_cpu.cpp
View 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);
}

1
intern/cycles/kernel/CMakeLists.txt
View 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

18
intern/cycles/kernel/geom/geom_object.h
View 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)

8
intern/cycles/kernel/kernel_globals.h
View 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;

94
intern/cycles/kernel/kernel_id_passes.h Normal file
View File

@ -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

93
intern/cycles/kernel/kernel_montecarlo.h
View File

@ -187,7 +187,10 @@ ccl_device float2 regular_polygon_sample(float corners, float rotation, float u,
ccl_device float3 ensure_valid_reflection(float3 Ng, float3 I, float3 N)
{
float3 R = 2*dot(N, I)*N - I;
if(dot(Ng, R) >= 0.05f) {
/* Reflection rays may always be at least as shallow as the incoming ray. */
float threshold = min(0.9f*dot(Ng, I), 0.01f);
if(dot(Ng, R) >= threshold) {
return N;
}
@ -195,22 +198,86 @@ ccl_device float3 ensure_valid_reflection(float3 Ng, float3 I, float3 N)
* The X axis is found by normalizing the component of N that's orthogonal to Ng.
* The Y axis isn't actually needed.
*/
float3 X = normalize(N - dot(N, Ng)*Ng);
float NdotNg = dot(N, Ng);
float3 X = normalize(N - NdotNg*Ng);
/* Calculate N.z and N.x in the local coordinate system. */
float Iz = dot(I, Ng);
float Ix2 = sqr(dot(I, X)), Iz2 = sqr(Iz);
float Ix2Iz2 = Ix2 + Iz2;
/* Calculate N.z and N.x in the local coordinate system.
*
* The goal of this computation is to find a N' that is rotated towards Ng just enough
* to lift R' above the threshold (here called t), therefore dot(R', Ng) = t.
*
* According to the standard reflection equation, this means that we want dot(2*dot(N', I)*N' - I, Ng) = t.
*
* Since the Z axis of our local coordinate system is Ng, dot(x, Ng) is just x.z, so we get 2*dot(N', I)*N'.z - I.z = t.
*
* The rotation is simple to express in the coordinate system we formed - since N lies in the X-Z-plane, we know that
* N' will also lie in the X-Z-plane, so N'.y = 0 and therefore dot(N', I) = N'.x*I.x + N'.z*I.z .
*
* Furthermore, we want N' to be normalized, so N'.x = sqrt(1 - N'.z^2).
*
* With these simplifications, we get the final equation 2*(sqrt(1 - N'.z^2)*I.x + N'.z*I.z)*N'.z - I.z = t.
*
* The only unknown here is N'.z, so we can solve for that.
*
* The equation has four solutions in general:
*
* N'.z = +-sqrt(0.5*(+-sqrt(I.x^2*(I.x^2 + I.z^2 - t^2)) + t*I.z + I.x^2 + I.z^2)/(I.x^2 + I.z^2))
* We can simplify this expression a bit by grouping terms:
*
* a = I.x^2 + I.z^2
* b = sqrt(I.x^2 * (a - t^2))
* c = I.z*t + a
* N'.z = +-sqrt(0.5*(+-b + c)/a)
*
* Two solutions can immediately be discarded because they're negative so N' would lie in the lower hemisphere.
*/
float Ix = dot(I, X), Iz = dot(I, Ng);
float Ix2 = sqr(Ix), Iz2 = sqr(Iz);
float a = Ix2 + Iz2;
float a = safe_sqrtf(Ix2*(Ix2Iz2 - sqr(0.05f)));
float b = Iz*0.05f + Ix2Iz2;
float c = (a + b > 0.0f)? (a + b) : (-a + b);
float b = safe_sqrtf(Ix2*(a - sqr(threshold)));
float c = Iz*threshold + a;
float Nz = safe_sqrtf(0.5f * c * (1.0f / Ix2Iz2));
float Nx = safe_sqrtf(1.0f - sqr(Nz));
/* Evaluate both solutions.
* In many cases one can be immediately discarded (if N'.z would be imaginary or larger than one), so check for that first.
* If no option is viable (might happen in extreme cases like N being in the wrong hemisphere), give up and return Ng. */
float fac = 0.5f/a;
float N1_z2 = fac*(b+c), N2_z2 = fac*(-b+c);
bool valid1 = (N1_z2 > 1e-5f) && (N1_z2 <= (1.0f + 1e-5f));
bool valid2 = (N2_z2 > 1e-5f) && (N2_z2 <= (1.0f + 1e-5f));
/* Transform back into global coordinates. */
return Nx*X + Nz*Ng;
float2 N_new;
if(valid1 && valid2) {
/* If both are possible, do the expensive reflection-based check. */
float2 N1 = make_float2(safe_sqrtf(1.0f - N1_z2), safe_sqrtf(N1_z2));
float2 N2 = make_float2(safe_sqrtf(1.0f - N2_z2), safe_sqrtf(N2_z2));
float R1 = 2*(N1.x*Ix + N1.y*Iz)*N1.y - Iz;
float R2 = 2*(N2.x*Ix + N2.y*Iz)*N2.y - Iz;
valid1 = (R1 >= 1e-5f);
valid2 = (R2 >= 1e-5f);
if(valid1 && valid2) {
/* If both solutions are valid, return the one with the shallower reflection since it will be closer to the input
* (if the original reflection wasn't shallow, we would not be in this part of the function). */
N_new = (R1 < R2)? N1 : N2;
}
else {
/* If only one reflection is valid (= positive), pick that one. */
N_new = (R1 > R2)? N1 : N2;
}
}
else if(valid1 || valid2) {
/* Only one solution passes the N'.z criterium, so pick that one. */
float Nz2 = valid1? N1_z2 : N2_z2;
N_new = make_float2(safe_sqrtf(1.0f - Nz2), safe_sqrtf(Nz2));
}
else {
return Ng;
}
return N_new.x*X + N_new.y*Ng;
}
CCL_NAMESPACE_END

43
intern/cycles/kernel/kernel_passes.h
View File

@ -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))

5
intern/cycles/kernel/kernel_shader.h
View 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

23
intern/cycles/kernel/kernel_types.h
View 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,6 +1270,9 @@ 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;
@ -1270,7 +1283,7 @@ typedef struct KernelFilm {
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;

26
intern/cycles/kernel/kernels/cuda/kernel.cu
View 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) {
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) {
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

12
intern/cycles/kernel/kernels/opencl/kernel.cl
View File

@ -66,10 +66,18 @@ __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__ */

26
intern/cycles/kernel/osl/osl_services.cpp
View File

@ -884,6 +884,23 @@ bool OSLRenderServices::has_userdata(ustring name, TypeDesc type, OSL::ShaderGlo
return false; /* never called by OSL */
}
TextureSystem::TextureHandle *OSLRenderServices::get_texture_handle(ustring filename)
{
if (filename.length() && filename[0] == '@') {
/* Dummy, we don't use texture handles for builtin textures but need
* to tell the OSL runtime optimizer that this is a valid texture. */
return NULL;
}
else {
return texturesys()->get_texture_handle(filename);
}
}
bool OSLRenderServices::good(TextureSystem::TextureHandle *texture_handle)
{
return texturesys()->good(texture_handle);
}
bool OSLRenderServices::texture(ustring filename,
TextureHandle *texture_handle,
TexturePerthread *texture_thread_info,
@ -894,7 +911,8 @@ bool OSLRenderServices::texture(ustring filename,
int nchannels,
float *result,
float *dresultds,
float *dresultdt)
float *dresultdt,
ustring *errormessage)
{
OSL::TextureSystem *ts = osl_ts;
ShaderData *sd = (ShaderData *)(sg->renderstate);
@ -1156,8 +1174,14 @@ bool OSLRenderServices::get_texture_info(OSL::ShaderGlobals *sg, ustring filenam
TypeDesc datatype, void *data)
{
OSL::TextureSystem *ts = osl_ts;
if (filename.length() && filename[0] == '@') {
/* Special builtin textures. */
return false;
}
else {
return ts->get_texture_info(filename, subimage, dataname, datatype, data);
}
}
int OSLRenderServices::pointcloud_search(OSL::ShaderGlobals *sg, ustring filename, const OSL::Vec3 &center,
float radius, int max_points, bool sort,

7
intern/cycles/kernel/osl/osl_services.h
View File

@ -93,6 +93,10 @@ public:
bool getmessage(OSL::ShaderGlobals *sg, ustring source, ustring name,
TypeDesc type, void *val, bool derivatives);
TextureSystem::TextureHandle *get_texture_handle(ustring filename);
bool good(TextureSystem::TextureHandle *texture_handle);
bool texture(ustring filename,
TextureSystem::TextureHandle *texture_handle,
TexturePerthread *texture_thread_info,
@ -103,7 +107,8 @@ public:
int nchannels,
float *result,
float *dresultds,
float *dresultdt);
float *dresultdt,
ustring *errormessage);
bool texture3d(ustring filename,
TextureHandle *texture_handle,

62
intern/cycles/kernel/shaders/stdosl.h
View File

@ -284,33 +284,63 @@ point rotate (point p, float angle, point a, point b)
normal ensure_valid_reflection(normal Ng, vector I, normal N)
{
/* The implementation here mirrors the one in kernel_montecarlo.h,
* check there for an explanation of the algorithm. */
float sqr(float x) { return x*x; }
vector R = 2*dot(N, I)*N - I;
if (dot(Ng, R) >= 0.05) {
float threshold = min(0.9*dot(Ng, I), 0.01);
if(dot(Ng, R) >= threshold) {
return N;
}
/* Form coordinate system with Ng as the Z axis and N inside the X-Z-plane.
* The X axis is found by normalizing the component of N that's orthogonal to Ng.
* The Y axis isn't actually needed.
*/
vector X = normalize(N - dot(N, Ng)*Ng);
float NdotNg = dot(N, Ng);
vector X = normalize(N - NdotNg*Ng);
/* Calculate N.z and N.x in the local coordinate system. */
float Ix = dot(I, X), Iz = dot(I, Ng);
float Ix2 = sqr(dot(I, X)), Iz2 = sqr(dot(I, Ng));
float Ix2Iz2 = Ix2 + Iz2;
float Ix2 = sqr(Ix), Iz2 = sqr(Iz);
float a = Ix2 + Iz2;
float a = sqrt(Ix2*(Ix2Iz2 - sqr(0.05)));
float b = Iz*0.05 + Ix2Iz2;
float c = (a + b > 0.0)? (a + b) : (-a + b);
float b = sqrt(Ix2*(a - sqr(threshold)));
float c = Iz*threshold + a;
float Nz = sqrt(0.5 * c * (1.0 / Ix2Iz2));
float Nx = sqrt(1.0 - sqr(Nz));
float fac = 0.5/a;
float N1_z2 = fac*(b+c), N2_z2 = fac*(-b+c);
int valid1 = (N1_z2 > 1e-5) && (N1_z2 <= (1.0 + 1e-5));
int valid2 = (N2_z2 > 1e-5) && (N2_z2 <= (1.0 + 1e-5));
/* Transform back into global coordinates. */
return Nx*X + Nz*Ng;
float N_new_x, N_new_z;
if(valid1 && valid2) {
float N1_x = sqrt(1.0 - N1_z2), N1_z = sqrt(N1_z2);
float N2_x = sqrt(1.0 - N2_z2), N2_z = sqrt(N2_z2);
float R1 = 2*(N1_x*Ix + N1_z*Iz)*N1_z - Iz;
float R2 = 2*(N2_x*Ix + N2_z*Iz)*N2_z - Iz;
valid1 = (R1 >= 1e-5);
valid2 = (R2 >= 1e-5);
if(valid1 && valid2) {
N_new_x = (R1 < R2)? N1_x : N2_x;
N_new_z = (R1 < R2)? N1_z : N2_z;
}
else {
N_new_x = (R1 > R2)? N1_x : N2_x;
N_new_z = (R1 > R2)? N1_z : N2_z;
}
}
else if(valid1 || valid2) {
float Nz2 = valid1? N1_z2 : N2_z2;
N_new_x = sqrt(1.0 - Nz2);
N_new_z = sqrt(Nz2);
}
else {
return Ng;
}
return N_new_x*X + N_new_z*Ng;
}

13
intern/cycles/kernel/split/kernel_buffer_update.h
View 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;

2
intern/cycles/render/CMakeLists.txt
View 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

56
intern/cycles/render/buffers.cpp
View File

@ -160,11 +160,12 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
(offset == DENOISING_PASS_DEPTH_VAR) ||
(offset == DENOISING_PASS_COLOR_VAR);
float scale_exposure = scale;
if(offset == DENOISING_PASS_COLOR || offset == DENOISING_PASS_CLEAN) {
scale *= exposure;
scale_exposure *= exposure;
}
else if(offset == DENOISING_PASS_COLOR_VAR) {
scale *= exposure*exposure;
scale_exposure *= exposure*exposure;
}
offset += params.get_denoising_offset();
@ -181,14 +182,14 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
if(components == 1) {
for(int i = 0; i < size; i++, mean += pass_stride, var += pass_stride, pixels++) {
pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale;
pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale_exposure;
}
}
else if(components == 3) {
for(int i = 0; i < size; i++, mean += pass_stride, var += pass_stride, pixels += 3) {
pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale;
pixels[1] = max(0.0f, var[1] - mean[1]*mean[1]*invsample)*scale;
pixels[2] = max(0.0f, var[2] - mean[2]*mean[2]*invsample)*scale;
pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale_exposure;
pixels[1] = max(0.0f, var[1] - mean[1]*mean[1]*invsample)*scale_exposure;
pixels[2] = max(0.0f, var[2] - mean[2]*mean[2]*invsample)*scale_exposure;
}
}
else {
@ -200,14 +201,28 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
if(components == 1) {
for(int i = 0; i < size; i++, in += pass_stride, pixels++) {
pixels[0] = in[0]*scale;
pixels[0] = in[0]*scale_exposure;
}
}
else if(components == 3) {
for(int i = 0; i < size; i++, in += pass_stride, pixels += 3) {
pixels[0] = in[0]*scale;
pixels[1] = in[1]*scale;
pixels[2] = in[2]*scale;
pixels[0] = in[0]*scale_exposure;
pixels[1] = in[1]*scale_exposure;
pixels[2] = in[2]*scale_exposure;
}
}
else if(components == 4) {
assert(offset == DENOISING_PASS_COLOR);
/* Since the alpha channel is not involved in denoising, output the Combined alpha channel. */
assert(params.passes[0].type == PASS_COMBINED);
float *in_combined = buffer.data();
for(int i = 0; i < size; i++, in += pass_stride, in_combined += pass_stride, pixels += 4) {
pixels[0] = in[0]*scale_exposure;
pixels[1] = in[1]*scale_exposure;
pixels[2] = in[2]*scale_exposure;
pixels[3] = saturate(in_combined[3]*scale);
}
}
else {
@ -218,7 +233,7 @@ bool RenderBuffers::get_denoising_pass_rect(int offset, float exposure, int samp
return true;
}
bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels)
bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int components, float *pixels, const string &name)
{
if(buffer.data() == NULL) {
return false;
@ -234,6 +249,14 @@ bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int
continue;
}
/* Tell Cryptomatte passes apart by their name. */
if(pass.type == PASS_CRYPTOMATTE) {
if(pass.name != name) {
pass_offset += pass.components;
continue;
}
}
float *in = buffer.data() + pass_offset;
int pass_stride = params.get_passes_size();
@ -370,6 +393,17 @@ bool RenderBuffers::get_pass_rect(PassType type, float exposure, int sample, int
pixels[3] = f.w*invw;
}
}
else if(type == PASS_CRYPTOMATTE) {
for(int i = 0; i < size; i++, in += pass_stride, pixels += 4) {
float4 f = make_float4(in[0], in[1], in[2], in[3]);
/* x and z contain integer IDs, don't rescale them.
y and w contain matte weights, they get scaled. */
pixels[0] = f.x;
pixels[1] = f.y * scale;
pixels[2] = f.z;
pixels[3] = f.w * scale;
}
}
else {
for(int i = 0; i < size; i++, in += pass_stride, pixels += 4) {
float4 f = make_float4(in[0], in[1], in[2], in[3]);

4
intern/cycles/render/buffers.h
View 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);
};

143
intern/cycles/render/coverage.cpp Normal file
View File

@ -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

49
intern/cycles/render/coverage.h Normal file
View File

@ -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__ */

37
intern/cycles/render/film.cpp
View 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);

13
intern/cycles/render/film.h
View 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);
};

5
intern/cycles/render/object.cpp
View 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) {

1
intern/cycles/render/object.h
View 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;

4
intern/cycles/render/shader.cpp
View 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>
@ -524,12 +525,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;

2
intern/cycles/util/CMakeLists.txt
View 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

28
intern/cycles/util/util_atomic.h
View 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()

127
intern/cycles/util/util_murmurhash.cpp Normal file
View File

@ -0,0 +1,127 @@
/*
* Copyright 2018 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
/* This is taken from alShaders/Cryptomatte/MurmurHash3.h:
*
* MurmurHash3 was written by Austin Appleby, and is placed in the public
* domain. The author hereby disclaims copyright to this source code.
*
*/
#include <stdlib.h>
#include <string.h>
#include "util/util_algorithm.h"
#include "util/util_murmurhash.h"
#if defined(_MSC_VER)
# define ROTL32(x,y) _rotl(x,y)
# define ROTL64(x,y) _rotl64(x,y)
# define BIG_CONSTANT(x) (x)
#else
ccl_device_inline uint32_t rotl32(uint32_t x, int8_t r)
{
return (x << r) | (x >> (32 - r));
}
# define ROTL32(x,y) rotl32(x,y)
# define BIG_CONSTANT(x) (x##LLU)
#endif
CCL_NAMESPACE_BEGIN
/* Block read - if your platform needs to do endian-swapping or can only
* handle aligned reads, do the conversion here. */
ccl_device_inline uint32_t mm_hash_getblock32(const uint32_t *p, int i)
{
return p[i];
}
/* Finalization mix - force all bits of a hash block to avalanche */
ccl_device_inline uint32_t mm_hash_fmix32 ( uint32_t h )
{
h ^= h >> 16;
h *= 0x85ebca6b;
h ^= h >> 13;
h *= 0xc2b2ae35;
h ^= h >> 16;
return h;
}
uint32_t util_murmur_hash3(const void *key, int len, uint32_t seed)
{
const uint8_t * data = (const uint8_t*)key;
const int nblocks = len / 4;
uint32_t h1 = seed;
const uint32_t c1 = 0xcc9e2d51;
const uint32_t c2 = 0x1b873593;
const uint32_t * blocks = (const uint32_t *)(data + nblocks*4);
for(int i = -nblocks; i; i++) {
uint32_t k1 = mm_hash_getblock32(blocks,i);
k1 *= c1;
k1 = ROTL32(k1,15);
k1 *= c2;
h1 ^= k1;
h1 = ROTL32(h1,13);
h1 = h1 * 5 + 0xe6546b64;
}
const uint8_t *tail = (const uint8_t*)(data + nblocks*4);
uint32_t k1 = 0;
switch(len & 3) {
case 3:
k1 ^= tail[2] << 16;
ATTR_FALLTHROUGH;
case 2:
k1 ^= tail[1] << 8;
ATTR_FALLTHROUGH;
case 1:
k1 ^= tail[0];
k1 *= c1;
k1 = ROTL32(k1,15);
k1 *= c2;
h1 ^= k1;
}
h1 ^= len;
h1 = mm_hash_fmix32(h1);
return h1;
}
/* This is taken from the cryptomatte specification 1.0 */
float util_hash_to_float(uint32_t hash)
{
uint32_t mantissa = hash & (( 1 << 23) - 1);
uint32_t exponent = (hash >> 23) & ((1 << 8) - 1);
exponent = max(exponent, (uint32_t) 1);
exponent = min(exponent, (uint32_t) 254);
exponent = exponent << 23;
uint32_t sign = (hash >> 31);
sign = sign << 31;
uint32_t float_bits = sign | exponent | mantissa;
float f;
memcpy(&f, &float_bits, sizeof(uint32_t));
return f;
}
CCL_NAMESPACE_END

30
intern/cycles/util/util_murmurhash.h Normal file
View File

@ -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__ */

2
release/scripts/addons

@ -1 +1 @@
Subproject commit 2e14c2aa69726b472f8758f62ad839eddaf63bfe
Subproject commit 6c3a46dc113de870a03191e4c0685238b0823acd

2
release/scripts/addons_contrib

@ -1 +1 @@
Subproject commit 311b03bd2ce4c5b3c3fc5e2b58a4ee1a629ea6a9
Subproject commit 15b25a42783d1e516b5298d70b582fae2559ae17

2
source/blender/editors/space_file/fsmenu.c
View File

@ -603,6 +603,8 @@ void fsmenu_read_system(struct FSMenu *fsmenu, int read_bookmarks)
/* not sure if this is right, but seems to give the relevant mnts */
if (!STREQLEN(mnt->mnt_fsname, "/dev", 4))
continue;
if (STREQLEN(mnt->mnt_fsname, "/dev/loop", 9))
continue;
len = strlen(mnt->mnt_dir);
if (len && mnt->mnt_dir[len - 1] != '/') {

2
source/tools

@ -1 +1 @@
Subproject commit 5162393c104d6d5f0314183a084875fff68f28bb
Subproject commit 11656ebaf7f912cdb1b5eb39c5d0a3b5d492c1aa