blender
/
blender
133
398
Fork 579
Code Issues 6.2k Pull Requests 745 Projects 19 Wiki Activity

Mesh: Replace auto smooth with node group #108014

Merged
Hans Goudey merged 149 commits from HooglyBoogly/blender:refactor-mesh-corner-normals-lazy into main 2023-10-20 16:54:20 +02:00
Conversation 15 Commits 149 Files Changed 98 +5679 -1791
311 changed files with 5679 additions and 1791 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files
Showing only changes of commit 5d05503945 - Show all commits

3
CMakeLists.txt
View File

@ -445,6 +445,9 @@ endif()
option(WITH_PYTHON_INSTALL "Copy system python into the blender install folder" ON)
option(WITH_INSTALL_COPYRIGHT "Copy the official Blender Foundation's copyright.txt into the Blender install folder" OFF)
mark_as_advanced(WITH_INSTALL_COPYRIGHT)
if((WITH_AUDASPACE AND NOT WITH_SYSTEM_AUDASPACE) OR WITH_MOD_FLUID)
option(WITH_PYTHON_NUMPY "Include NumPy in Blender (used by Audaspace and Mantaflow)" ON)
endif()

5
doc/manpage/blender.1.py
View File

@ -31,7 +31,10 @@ def man_format(data: str) -> str:
def blender_extract_info(blender_bin: str) -> Dict[str, str]:
blender_env = {
"ASAN_OPTIONS": "exitcode=0:" + os.environ.get("ASAN_OPTIONS", ""),
"ASAN_OPTIONS": (
os.environ.get("ASAN_OPTIONS", "") +
":exitcode=0:check_initialization_order=0:strict_init_order=0"
).lstrip(":"),
}
blender_help = subprocess.run(

2
intern/cycles/blender/CMakeLists.txt
View File

@ -26,6 +26,7 @@ set(SRC
image.cpp
geometry.cpp
light.cpp
light_linking.cpp
mesh.cpp
object.cpp
object_cull.cpp
@ -47,6 +48,7 @@ set(SRC
display_driver.h
id_map.h
image.h
light_linking.h
object_cull.h
output_driver.h
sync.h

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

@ -7,7 +7,7 @@ from bpy.app.translations import contexts as i18n_contexts
from bpy_extras.node_utils import find_node_input
from bl_ui.utils import PresetPanel
from bpy.types import Panel
from bpy.types import Panel, Menu
from bl_ui.properties_grease_pencil_common import GreasePencilSimplifyPanel
from bl_ui.properties_render import draw_curves_settings
@ -1301,6 +1301,7 @@ class CYCLES_OBJECT_PT_lightgroup(CyclesButtonsPanel, Panel):
bl_label = "Light Group"
bl_parent_id = "CYCLES_OBJECT_PT_shading"
bl_context = "object"
bl_options = {'DEFAULT_CLOSED'}
def draw(self, context):
layout = self.layout
@ -1321,6 +1322,88 @@ class CYCLES_OBJECT_PT_lightgroup(CyclesButtonsPanel, Panel):
sub.operator("scene.view_layer_add_lightgroup", icon='ADD', text="").name = ob.lightgroup
class CYCLES_OBJECT_MT_light_linking_context_menu(Menu):
bl_label = "Light Linking Specials"
def draw(self, _context):
layout = self.layout
layout.operator("object.light_linking_receivers_select")
class CYCLES_OBJECT_MT_shadow_linking_context_menu(Menu):
bl_label = "Shadow Linking Specials"
def draw(self, _context):
layout = self.layout
layout.operator("object.light_linking_blockers_select")
class CYCLES_OBJECT_PT_light_linking(CyclesButtonsPanel, Panel):
bl_label = "Light Linking"
bl_parent_id = "CYCLES_OBJECT_PT_shading"
bl_context = "object"
bl_options = {'DEFAULT_CLOSED'}
def draw(self, context):
layout = self.layout
layout.use_property_split = True
object = context.object
light_linking = object.light_linking
col = layout.column()
col.template_ID(
light_linking,
"receiver_collection",
new="object.light_linking_receiver_collection_new")
if not light_linking.receiver_collection:
return
row = layout.row()
col = row.column()
col.template_light_linking_collection(light_linking, "receiver_collection")
col = row.column()
sub = col.column(align=True)
sub.menu("CYCLES_OBJECT_MT_light_linking_context_menu", icon='DOWNARROW_HLT', text="")
class CYCLES_OBJECT_PT_shadow_linking(CyclesButtonsPanel, Panel):
bl_label = "Shadow Linking"
bl_parent_id = "CYCLES_OBJECT_PT_shading"
bl_context = "object"
bl_options = {'DEFAULT_CLOSED'}
def draw(self, context):
layout = self.layout
layout.use_property_split = True
object = context.object
light_linking = object.light_linking
col = layout.column()
col.template_ID(
light_linking,
"blocker_collection",
new="object.light_linking_blocker_collection_new")
if not light_linking.blocker_collection:
return
row = layout.row()
col = row.column()
col.template_light_linking_collection(light_linking, "blocker_collection")
col = row.column()
sub = col.column(align=True)
sub.menu("CYCLES_OBJECT_MT_shadow_linking_context_menu", icon='DOWNARROW_HLT', text="")
class CYCLES_OBJECT_PT_visibility(CyclesButtonsPanel, Panel):
bl_label = "Visibility"
bl_context = "object"
@ -2451,6 +2534,10 @@ classes = (
CYCLES_OBJECT_PT_shading_gi_approximation,
CYCLES_OBJECT_PT_shading_caustics,
CYCLES_OBJECT_PT_lightgroup,
CYCLES_OBJECT_MT_light_linking_context_menu,
CYCLES_OBJECT_PT_light_linking,
CYCLES_OBJECT_MT_shadow_linking_context_menu,
CYCLES_OBJECT_PT_shadow_linking,
CYCLES_OBJECT_PT_visibility,
CYCLES_OBJECT_PT_visibility_ray_visibility,
CYCLES_OBJECT_PT_visibility_culling,

7
intern/cycles/blender/light.cpp
View File

@ -3,6 +3,7 @@
#include "scene/light.h"
#include "blender/light_linking.h"
#include "blender/sync.h"
#include "blender/util.h"
@ -145,8 +146,12 @@ void BlenderSync::sync_light(BL::Object &b_parent,
light->set_use_scatter((visibility & PATH_RAY_VOLUME_SCATTER) != 0);
light->set_is_shadow_catcher(b_ob_info.real_object.is_shadow_catcher());
/* lightgroup */
/* Light group and linking. */
light->set_lightgroup(ustring(b_ob_info.real_object.lightgroup()));
light->set_light_set_membership(
BlenderLightLink::get_light_set_membership(PointerRNA_NULL, b_ob_info.real_object));
light->set_shadow_set_membership(
BlenderLightLink::get_shadow_set_membership(PointerRNA_NULL, b_ob_info.real_object));
/* tag */
light->tag_update(scene);

63
intern/cycles/blender/light_linking.cpp Normal file
View File

@ -0,0 +1,63 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2023 Blender Foundation */
#include "blender/light_linking.h"
#include "scene/object.h"
#include "DNA_object_types.h"
CCL_NAMESPACE_BEGIN
static const ::Object *get_blender_object(const BL::Object &object)
{
return reinterpret_cast<::Object *>(object.ptr.data);
}
static const ::LightLinking *get_light_linking(const BL::Object &object)
{
const ::Object *blender_object = get_blender_object(object);
return blender_object->light_linking;
}
uint64_t BlenderLightLink::get_light_set_membership(const BL::Object & /*parent*/,
const BL::Object &object)
{
const ::LightLinking *light_linking = get_light_linking(object);
return (light_linking) ? light_linking->runtime.light_set_membership : LIGHT_LINK_MASK_ALL;
}
uint BlenderLightLink::get_receiver_light_set(const BL::Object &parent, const BL::Object &object)
{
if (parent) {
const ::LightLinking *parent_light_linking = get_light_linking(parent);
if (parent_light_linking && parent_light_linking->runtime.receiver_light_set) {
return parent_light_linking->runtime.receiver_light_set;
}
}
const ::LightLinking *light_linking = get_light_linking(object);
return (light_linking) ? light_linking->runtime.receiver_light_set : 0;
}
uint64_t BlenderLightLink::get_shadow_set_membership(const BL::Object & /*parent*/,
const BL::Object &object)
{
const ::LightLinking *light_linking = get_light_linking(object);
return (light_linking) ? light_linking->runtime.shadow_set_membership : LIGHT_LINK_MASK_ALL;
}
uint BlenderLightLink::get_blocker_shadow_set(const BL::Object &parent, const BL::Object &object)
{
if (parent) {
const ::LightLinking *parent_light_linking = get_light_linking(parent);
if (parent_light_linking && parent_light_linking->runtime.blocker_shadow_set) {
return parent_light_linking->runtime.blocker_shadow_set;
}
}
const ::LightLinking *light_linking = get_light_linking(object);
return (light_linking) ? light_linking->runtime.blocker_shadow_set : 0;
}
CCL_NAMESPACE_END

22
intern/cycles/blender/light_linking.h Normal file
View File

@ -0,0 +1,22 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2023 Blender Foundation */
#pragma once
#include <cstdint>
#include "MEM_guardedalloc.h"
#include "RNA_blender_cpp.h"
CCL_NAMESPACE_BEGIN
class BlenderLightLink {
public:
static uint64_t get_light_set_membership(const BL::Object &parent, const BL::Object &object);
static uint get_receiver_light_set(const BL::Object &parent, const BL::Object &object);
static uint64_t get_shadow_set_membership(const BL::Object &parent, const BL::Object &object);
static uint get_blocker_shadow_set(const BL::Object &parent, const BL::Object &object);
};
CCL_NAMESPACE_END

8
intern/cycles/blender/object.cpp
View File

@ -1,6 +1,7 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2022 Blender Foundation */
#include "blender/light_linking.h"
#include "blender/object_cull.h"
#include "blender/sync.h"
#include "blender/util.h"
@ -348,9 +349,14 @@ Object *BlenderSync::sync_object(BL::Depsgraph &b_depsgraph,
object->set_random_id(hash_uint2(hash_string(object->name.c_str()), 0));
}
/* lightgroup */
/* Light group and linking. */
object->set_lightgroup(ustring(b_ob.lightgroup()));
object->set_light_set_membership(BlenderLightLink::get_light_set_membership(b_parent, b_ob));
object->set_receiver_light_set(BlenderLightLink::get_receiver_light_set(b_parent, b_ob));
object->set_shadow_set_membership(BlenderLightLink::get_shadow_set_membership(b_parent, b_ob));
object->set_blocker_shadow_set(BlenderLightLink::get_blocker_shadow_set(b_parent, b_ob));
object->tag_update(scene);
}

2
intern/cycles/device/cpu/kernel.cpp
View File

@ -24,11 +24,13 @@ CPUKernels::CPUKernels()
REGISTER_KERNEL(integrator_intersect_shadow),
REGISTER_KERNEL(integrator_intersect_subsurface),
REGISTER_KERNEL(integrator_intersect_volume_stack),
REGISTER_KERNEL(integrator_intersect_dedicated_light),
REGISTER_KERNEL(integrator_shade_background),
REGISTER_KERNEL(integrator_shade_light),
REGISTER_KERNEL(integrator_shade_shadow),
REGISTER_KERNEL(integrator_shade_surface),
REGISTER_KERNEL(integrator_shade_volume),
REGISTER_KERNEL(integrator_shade_dedicated_light),
REGISTER_KERNEL(integrator_megakernel),
/* Shader evaluation. */
REGISTER_KERNEL(shader_eval_displace),

2
intern/cycles/device/cpu/kernel.h
View File

@ -33,11 +33,13 @@ class CPUKernels {
IntegratorFunction integrator_intersect_shadow;
IntegratorFunction integrator_intersect_subsurface;
IntegratorFunction integrator_intersect_volume_stack;
IntegratorFunction integrator_intersect_dedicated_light;
IntegratorShadeFunction integrator_shade_background;
IntegratorShadeFunction integrator_shade_light;
IntegratorShadeFunction integrator_shade_shadow;
IntegratorShadeFunction integrator_shade_surface;
IntegratorShadeFunction integrator_shade_volume;
IntegratorShadeFunction integrator_shade_dedicated_light;
IntegratorShadeFunction integrator_megakernel;
/* Shader evaluation. */

6
intern/cycles/device/kernel.cpp
View File

@ -18,6 +18,7 @@ bool device_kernel_has_shading(DeviceKernel kernel)
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT ||
kernel == DEVICE_KERNEL_SHADER_EVAL_DISPLACE ||
kernel == DEVICE_KERNEL_SHADER_EVAL_BACKGROUND ||
kernel == DEVICE_KERNEL_SHADER_EVAL_CURVE_SHADOW_TRANSPARENCY);
@ -29,6 +30,7 @@ bool device_kernel_has_intersection(DeviceKernel kernel)
kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW ||
kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE ||
kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK ||
kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE);
}
@ -49,6 +51,8 @@ const char *device_kernel_as_string(DeviceKernel kernel)
return "integrator_intersect_subsurface";
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
return "integrator_intersect_volume_stack";
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
return "integrator_intersect_dedicated_light";
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
return "integrator_shade_background";
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
@ -63,6 +67,8 @@ const char *device_kernel_as_string(DeviceKernel kernel)
return "integrator_shade_surface_mnee";
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
return "integrator_shade_volume";
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
return "integrator_shade_dedicated_light";
case DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL:
return "integrator_megakernel";
case DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY:

4
intern/cycles/device/metal/kernel.mm
View File

@ -36,6 +36,10 @@ struct ShaderCache {
ShaderCache(id<MTLDevice> _mtlDevice) : mtlDevice(_mtlDevice)
{
/* Initialize occupancy tuning LUT. */
// TODO: Look into tuning for DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT and
// DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT.
if (MetalInfo::get_device_vendor(mtlDevice) == METAL_GPU_APPLE) {
switch (MetalInfo::get_apple_gpu_architecture(mtlDevice)) {
default:

1
intern/cycles/device/metal/queue.mm
View File

@ -531,6 +531,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
break;

10
intern/cycles/device/optix/device_impl.cpp
View File

@ -421,6 +421,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
group_descs[PG_RGEN_INTERSECT_VOLUME_STACK].raygen.module = optix_module;
group_descs[PG_RGEN_INTERSECT_VOLUME_STACK].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_intersect_volume_stack";
group_descs[PG_RGEN_INTERSECT_DEDICATED_LIGHT].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_INTERSECT_DEDICATED_LIGHT].raygen.module = optix_module;
group_descs[PG_RGEN_INTERSECT_DEDICATED_LIGHT].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_intersect_dedicated_light";
group_descs[PG_MISS].kind = OPTIX_PROGRAM_GROUP_KIND_MISS;
group_descs[PG_MISS].miss.module = optix_module;
group_descs[PG_MISS].miss.entryFunctionName = "__miss__kernel_optix_miss";
@ -547,6 +551,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
group_descs[PG_RGEN_SHADE_SHADOW].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_SHADOW].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_shade_shadow";
group_descs[PG_RGEN_SHADE_DEDICATED_LIGHT].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_SHADE_DEDICATED_LIGHT].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_DEDICATED_LIGHT].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_shade_dedicated_light";
group_descs[PG_RGEN_EVAL_DISPLACE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_EVAL_DISPLACE].raygen.module = optix_module;
group_descs[PG_RGEN_EVAL_DISPLACE].raygen.entryFunctionName =
@ -659,6 +667,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SHADOW]);
pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_SUBSURFACE]);
pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_VOLUME_STACK]);
pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_DEDICATED_LIGHT]);
pipeline_groups.push_back(groups[PG_MISS]);
pipeline_groups.push_back(groups[PG_HITD]);
pipeline_groups.push_back(groups[PG_HITS]);
@ -948,6 +957,7 @@ bool OptiXDevice::load_osl_kernels()
pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_MNEE]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_VOLUME]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_SHADOW]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_DEDICATED_LIGHT]);
pipeline_groups.push_back(groups[PG_RGEN_EVAL_DISPLACE]);
pipeline_groups.push_back(groups[PG_RGEN_EVAL_BACKGROUND]);
pipeline_groups.push_back(groups[PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY]);

2
intern/cycles/device/optix/device_impl.h
View File

@ -21,6 +21,7 @@ enum {
PG_RGEN_INTERSECT_SHADOW,
PG_RGEN_INTERSECT_SUBSURFACE,
PG_RGEN_INTERSECT_VOLUME_STACK,
PG_RGEN_INTERSECT_DEDICATED_LIGHT,
PG_RGEN_SHADE_BACKGROUND,
PG_RGEN_SHADE_LIGHT,
PG_RGEN_SHADE_SURFACE,
@ -28,6 +29,7 @@ enum {
PG_RGEN_SHADE_SURFACE_MNEE,
PG_RGEN_SHADE_VOLUME,
PG_RGEN_SHADE_SHADOW,
PG_RGEN_SHADE_DEDICATED_LIGHT,
PG_RGEN_EVAL_DISPLACE,
PG_RGEN_EVAL_BACKGROUND,
PG_RGEN_EVAL_CURVE_SHADOW_TRANSPARENCY,

9
intern/cycles/device/optix/queue.cpp
View File

@ -123,6 +123,10 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel,
pipeline = optix_device->pipelines[PIP_SHADE];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SHADOW * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
pipeline = optix_device->pipelines[PIP_SHADE];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_DEDICATED_LIGHT * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST:
pipeline = optix_device->pipelines[PIP_INTERSECT];
@ -140,6 +144,11 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel,
pipeline = optix_device->pipelines[PIP_INTERSECT];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_INTERSECT_VOLUME_STACK * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
pipeline = optix_device->pipelines[PIP_INTERSECT];
sbt_params.raygenRecord = sbt_data_ptr +
PG_RGEN_INTERSECT_DEDICATED_LIGHT * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_SHADER_EVAL_DISPLACE:
pipeline = optix_device->pipelines[PIP_SHADE];

24
intern/cycles/graph/node.cpp
View File

@ -68,6 +68,12 @@ void Node::set(const SocketType &input, uint value)
set_if_different(input, value);
}
void Node::set(const SocketType &input, uint64_t value)
{
assert(input.type == SocketType::UINT64);
set_if_different(input, value);
}
void Node::set(const SocketType &input, float value)
{
assert(input.type == SocketType::FLOAT);
@ -190,6 +196,12 @@ uint Node::get_uint(const SocketType &input) const
return get_socket_value<uint>(this, input);
}
uint64_t Node::get_uint64(const SocketType &input) const
{
assert(input.type == SocketType::UINT64);
return get_socket_value<uint64_t>(this, input);
}
float Node::get_float(const SocketType &input) const
{
assert(input.type == SocketType::FLOAT);
@ -434,6 +446,9 @@ void Node::set_value(const SocketType &socket, const Node &other, const SocketTy
case SocketType::UINT:
set(socket, get_socket_value<uint>(&other, socket));
break;
case SocketType::UINT64:
set(socket, get_socket_value<uint64_t>(&other, socket));
break;
case SocketType::COLOR:
case SocketType::VECTOR:
case SocketType::POINT:
@ -489,6 +504,8 @@ bool Node::equals_value(const Node &other, const SocketType &socket) const
return is_value_equal<int>(this, &other, socket);
case SocketType::UINT:
return is_value_equal<uint>(this, &other, socket);
case SocketType::UINT64:
return is_value_equal<uint64_t>(this, &other, socket);
case SocketType::COLOR:
return is_value_equal<float3>(this, &other, socket);
case SocketType::VECTOR:
@ -534,6 +551,7 @@ bool Node::equals_value(const Node &other, const SocketType &socket) const
return is_array_equal<void *>(this, &other, socket);
case SocketType::UNDEFINED:
case SocketType::NUM_TYPES:
return true;
}
@ -608,6 +626,9 @@ void Node::hash(MD5Hash &md5)
case SocketType::UINT:
value_hash<uint>(this, socket, md5);
break;
case SocketType::UINT64:
value_hash<uint64_t>(this, socket, md5);
break;
case SocketType::COLOR:
float3_hash(this, socket, md5);
break;
@ -673,6 +694,7 @@ void Node::hash(MD5Hash &md5)
break;
case SocketType::UNDEFINED:
case SocketType::NUM_TYPES:
break;
}
}
@ -697,6 +719,7 @@ size_t Node::get_total_size_in_bytes() const
case SocketType::FLOAT:
case SocketType::INT:
case SocketType::UINT:
case SocketType::UINT64:
case SocketType::COLOR:
case SocketType::VECTOR:
case SocketType::POINT:
@ -745,6 +768,7 @@ size_t Node::get_total_size_in_bytes() const
break;
case SocketType::UNDEFINED:
case SocketType::NUM_TYPES:
break;
}
}

2
intern/cycles/graph/node.h
View File

@ -95,6 +95,7 @@ struct Node {
void set(const SocketType &input, bool value);
void set(const SocketType &input, int value);
void set(const SocketType &input, uint value);
void set(const SocketType &input, uint64_t value);
void set(const SocketType &input, float value);
void set(const SocketType &input, float2 value);
void set(const SocketType &input, float3 value);
@ -127,6 +128,7 @@ struct Node {
bool get_bool(const SocketType &input) const;
int get_int(const SocketType &input) const;
uint get_uint(const SocketType &input) const;
uint64_t get_uint64(const SocketType &input) const;
float get_float(const SocketType &input) const;
float2 get_float2(const SocketType &input) const;
float3 get_float3(const SocketType &input) const;

17
intern/cycles/graph/node_type.cpp
View File

@ -23,6 +23,7 @@ size_t SocketType::size(Type type)
{
switch (type) {
case UNDEFINED:
case NUM_TYPES:
return 0;
case BOOLEAN:
@ -33,6 +34,8 @@ size_t SocketType::size(Type type)
return sizeof(int);
case UINT:
return sizeof(uint);
case UINT64:
return sizeof(uint64_t);
case COLOR:
return sizeof(float3);
case VECTOR:
@ -99,11 +102,12 @@ ustring SocketType::type_name(Type type)
ustring("boolean"), ustring("float"),
ustring("int"), ustring("uint"),
ustring("color"), ustring("vector"),
ustring("point"), ustring("normal"),
ustring("point2"), ustring("closure"),
ustring("string"), ustring("enum"),
ustring("transform"), ustring("node"),
ustring("uint64"), ustring("color"),
ustring("vector"), ustring("point"),
ustring("normal"), ustring("point2"),
ustring("closure"), ustring("string"),
ustring("enum"), ustring("transform"),
ustring("node"),
ustring("array_boolean"), ustring("array_float"),
ustring("array_int"), ustring("array_color"),
@ -112,6 +116,9 @@ ustring SocketType::type_name(Type type)
ustring("array_string"), ustring("array_transform"),
ustring("array_node")};
constexpr size_t num_names = sizeof(names) / sizeof(*names);
static_assert(num_names == NUM_TYPES);
return names[(int)type];
}

7
intern/cycles/graph/node_type.h
View File

@ -21,12 +21,13 @@ typedef uint64_t SocketModifiedFlags;
struct SocketType {
enum Type {
UNDEFINED,
UNDEFINED = 0,
BOOLEAN,
FLOAT,
INT,
UINT,
UINT64,
COLOR,
VECTOR,
POINT,
@ -49,6 +50,8 @@ struct SocketType {
STRING_ARRAY,
TRANSFORM_ARRAY,
NODE_ARRAY,
NUM_TYPES,
};
enum Flags {
@ -191,6 +194,8 @@ struct NodeType {
SOCKET_DEFINE(name, ui_name, default_value, int, SocketType::INT, 0, ##__VA_ARGS__)
#define SOCKET_UINT(name, ui_name, default_value, ...) \
SOCKET_DEFINE(name, ui_name, default_value, uint, SocketType::UINT, 0, ##__VA_ARGS__)
#define SOCKET_UINT64(name, ui_name, default_value, ...) \
SOCKET_DEFINE(name, ui_name, default_value, uint64_t, SocketType::UINT64, 0, ##__VA_ARGS__)
#define SOCKET_FLOAT(name, ui_name, default_value, ...) \
SOCKET_DEFINE(name, ui_name, default_value, float, SocketType::FLOAT, 0, ##__VA_ARGS__)
#define SOCKET_COLOR(name, ui_name, default_value, ...) \

10
intern/cycles/graph/node_xml.cpp
View File

@ -93,6 +93,10 @@ void xml_read_node(XMLReader &reader, Node *node, xml_node xml_node)
node->set(socket, (uint)atoi(attr.value()));
break;
}
case SocketType::UINT64: {
node->set(socket, (uint64_t)strtoull(attr.value(), nullptr, 10));
break;
}
case SocketType::INT_ARRAY: {
vector<string> tokens;
string_split(tokens, attr.value());
@ -213,6 +217,7 @@ void xml_read_node(XMLReader &reader, Node *node, xml_node xml_node)
}
case SocketType::CLOSURE:
case SocketType::UNDEFINED:
case SocketType::NUM_TYPES:
break;
}
}
@ -280,6 +285,10 @@ xml_node xml_write_node(Node *node, xml_node xml_root)
attr = node->get_uint(socket);
break;
}
case SocketType::UINT64: {
attr = node->get_uint64(socket);
break;
}
case SocketType::INT_ARRAY: {
std::stringstream ss;
const array<int> &value = node->get_int_array(socket);
@ -409,6 +418,7 @@ xml_node xml_write_node(Node *node, xml_node xml_root)
}
case SocketType::CLOSURE:
case SocketType::UNDEFINED:
case SocketType::NUM_TYPES:
break;
}
}

9
intern/cycles/integrator/path_trace_work_gpu.cpp
View File

@ -508,7 +508,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK: {
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT: {
/* Ray intersection kernels with integrator state. */
DeviceKernelArguments args(&d_path_index, &work_size);
@ -521,7 +522,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME: {
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT: {
/* Shading kernels with integrator state and render buffer. */
DeviceKernelArguments args(&d_path_index, &buffers_->buffer.device_pointer, &work_size);
@ -1177,7 +1179,8 @@ bool PathTraceWorkGPU::kernel_creates_shadow_paths(DeviceKernel kernel)
return (kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT);
}
bool PathTraceWorkGPU::kernel_creates_ao_paths(DeviceKernel kernel)

3
intern/cycles/kernel/CMakeLists.txt
View File

@ -267,6 +267,7 @@ set(SRC_KERNEL_INTEGRATOR_HEADERS
integrator/displacement_shader.h
integrator/init_from_bake.h
integrator/init_from_camera.h
integrator/intersect_dedicated_light.h
integrator/intersect_closest.h
integrator/intersect_shadow.h
integrator/intersect_subsurface.h
@ -280,7 +281,9 @@ set(SRC_KERNEL_INTEGRATOR_HEADERS
integrator/shade_shadow.h
integrator/shade_surface.h
integrator/shade_volume.h
integrator/shade_dedicated_light.h
integrator/shadow_catcher.h
integrator/shadow_linking.h
integrator/shadow_state_template.h
integrator/state_flow.h
integrator/state.h

6
intern/cycles/kernel/bvh/shadow_all.h
View File

@ -144,6 +144,12 @@ ccl_device_inline
continue;
}
#ifdef __SHADOW_LINKING__
if (intersection_skip_shadow_link(kg, ray, prim_object)) {
continue;
}
#endif
switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: {
hit = triangle_intersect(

6
intern/cycles/kernel/bvh/traversal.h
View File

@ -129,6 +129,12 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
continue;
}
#ifdef __SHADOW_LINKING__
if (intersection_skip_shadow_link(kg, ray, prim_object)) {
continue;
}
#endif
switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: {
if (triangle_intersect(kg,

37
intern/cycles/kernel/bvh/util.h
View File

@ -253,4 +253,41 @@ ccl_device_inline bool intersection_skip_self_local(ccl_private const RaySelfPri
return (self.prim == prim);
}
#ifdef __SHADOW_LINKING__
ccl_device_inline uint64_t ray_get_shadow_set_membership(KernelGlobals kg,
ccl_private const Ray *ray)
{
if (ray->self.light != LAMP_NONE) {
return kernel_data_fetch(lights, ray->self.light).shadow_set_membership;
}
if (ray->self.light_object != OBJECT_NONE) {
return kernel_data_fetch(objects, ray->self.light_object).shadow_set_membership;
}
return LIGHT_LINK_MASK_ALL;
}
#endif
ccl_device_inline bool intersection_skip_shadow_link(KernelGlobals kg,
ccl_private const Ray *ray,
const int isect_object)
{
#ifdef __SHADOW_LINKING__
if (!(kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING)) {
return false;
}
const uint64_t set_membership = ray_get_shadow_set_membership(kg, ray);
if (set_membership == LIGHT_LINK_MASK_ALL) {
return false;
}
const uint blocker_set = kernel_data_fetch(objects, isect_object).blocker_shadow_set;
return ((uint64_t(1) << uint64_t(blocker_set)) & set_membership) == 0;
#else
return false;
#endif
}
CCL_NAMESPACE_END

26
intern/cycles/kernel/device/cpu/bvh.h
View File

@ -178,14 +178,19 @@ ccl_device_inline void kernel_embree_setup_rayhit(const Ray &ray,
rayhit.hit.instID[0] = RTC_INVALID_GEOMETRY_ID;
}
ccl_device_inline int kernel_embree_get_hit_object(const RTCHit *hit)
{
return (hit->instID[0] != RTC_INVALID_GEOMETRY_ID ? hit->instID[0] : hit->geomID) / 2;
}
ccl_device_inline bool kernel_embree_is_self_intersection(const KernelGlobals kg,
const RTCHit *hit,
const Ray *ray,
const intptr_t prim_offset)
{
int object, prim;
object = (hit->instID[0] != RTC_INVALID_GEOMETRY_ID ? hit->instID[0] : hit->geomID) / 2;
const int object = kernel_embree_get_hit_object(hit);
int prim;
if ((ray->self.object == object) || (ray->self.light_object == object)) {
prim = hit->primID + prim_offset;
}
@ -209,7 +214,7 @@ ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg,
{
isect->t = ray->tfar;
isect->prim = hit->primID + prim_offset;
isect->object = hit->instID[0] != RTC_INVALID_GEOMETRY_ID ? hit->instID[0] / 2 : hit->geomID / 2;
isect->object = kernel_embree_get_hit_object(hit);
const bool is_hair = hit->geomID & 1;
if (is_hair) {
@ -287,7 +292,15 @@ ccl_device_forceinline void kernel_embree_filter_intersection_func_impl(
kg, hit, cray, reinterpret_cast<intptr_t>(args->geometryUserPtr)))
{
*args->valid = 0;
return;
}
#ifdef __SHADOW_LINKING__
if (intersection_skip_shadow_link(kg, cray, kernel_embree_get_hit_object(hit))) {
*args->valid = 0;
return;
}
#endif
}
/* This gets called by Embree at every valid ray/object intersection.
@ -323,6 +336,13 @@ ccl_device_forceinline void kernel_embree_filter_occluded_shadow_all_func_impl(
return;
}
#ifdef __SHADOW_LINKING__
if (intersection_skip_shadow_link(kg, cray, current_isect.object)) {
*args->valid = 0;
return;
}
#endif
/* If no transparent shadows or max number of hits exceeded, all light is blocked. */
const int flags = intersection_get_shader_flags(kg, current_isect.prim, current_isect.type);
if (!(flags & (SD_HAS_TRANSPARENT_SHADOW)) || ctx->num_hits >= ctx->max_hits) {

2
intern/cycles/kernel/device/cpu/kernel_arch.h
View File

@ -28,11 +28,13 @@ KERNEL_INTEGRATOR_SHADE_FUNCTION(intersect_closest);
KERNEL_INTEGRATOR_FUNCTION(intersect_shadow);
KERNEL_INTEGRATOR_FUNCTION(intersect_subsurface);
KERNEL_INTEGRATOR_FUNCTION(intersect_volume_stack);
KERNEL_INTEGRATOR_FUNCTION(intersect_dedicated_light);
KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_background);
KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_light);
KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_shadow);
KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_surface);
KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_volume);
KERNEL_INTEGRATOR_SHADE_FUNCTION(shade_dedicated_light);
KERNEL_INTEGRATOR_SHADE_FUNCTION(megakernel);
#undef KERNEL_INTEGRATOR_FUNCTION

2
intern/cycles/kernel/device/cpu/kernel_arch_impl.h
View File

@ -102,10 +102,12 @@ DEFINE_INTEGRATOR_INIT_KERNEL(init_from_bake)
DEFINE_INTEGRATOR_SHADE_KERNEL(intersect_closest)
DEFINE_INTEGRATOR_KERNEL(intersect_subsurface)
DEFINE_INTEGRATOR_KERNEL(intersect_volume_stack)
DEFINE_INTEGRATOR_KERNEL(intersect_dedicated_light)
DEFINE_INTEGRATOR_SHADE_KERNEL(shade_background)
DEFINE_INTEGRATOR_SHADE_KERNEL(shade_light)
DEFINE_INTEGRATOR_SHADE_KERNEL(shade_surface)
DEFINE_INTEGRATOR_SHADE_KERNEL(shade_volume)
DEFINE_INTEGRATOR_SHADE_KERNEL(shade_dedicated_light)
DEFINE_INTEGRATOR_SHADE_KERNEL(megakernel)
DEFINE_INTEGRATOR_SHADOW_KERNEL(intersect_shadow)
DEFINE_INTEGRATOR_SHADOW_SHADE_KERNEL(shade_shadow)

31
intern/cycles/kernel/device/gpu/kernel.h
View File

@ -27,10 +27,12 @@
#include "kernel/integrator/init_from_bake.h"
#include "kernel/integrator/init_from_camera.h"
#include "kernel/integrator/intersect_closest.h"
#include "kernel/integrator/intersect_dedicated_light.h"
#include "kernel/integrator/intersect_shadow.h"
#include "kernel/integrator/intersect_subsurface.h"
#include "kernel/integrator/intersect_volume_stack.h"
#include "kernel/integrator/shade_background.h"
#include "kernel/integrator/shade_dedicated_light.h"
#include "kernel/integrator/shade_light.h"
#include "kernel/integrator/shade_shadow.h"
#include "kernel/integrator/shade_surface.h"
@ -196,6 +198,20 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
}
ccl_gpu_kernel_postfix
ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
ccl_gpu_kernel_signature(integrator_intersect_dedicated_light,
ccl_global const int *path_index_array,
const int work_size)
{
const int global_index = ccl_gpu_global_id_x();
if (ccl_gpu_kernel_within_bounds(global_index, work_size)) {
const int state = (path_index_array) ? path_index_array[global_index] : global_index;
ccl_gpu_kernel_call(integrator_intersect_dedicated_light(NULL, state));
}
}
ccl_gpu_kernel_postfix
# ifdef __KERNEL_ONEAPI__
# include "kernel/device/oneapi/context_intersect_end.h"
# endif
@ -334,6 +350,21 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
}
ccl_gpu_kernel_postfix
ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
ccl_gpu_kernel_signature(integrator_shade_dedicated_light,
ccl_global const int *path_index_array,
ccl_global float *render_buffer,
const int work_size)
{
const int global_index = ccl_gpu_global_id_x();
if (ccl_gpu_kernel_within_bounds(global_index, work_size)) {
const int state = (path_index_array) ? path_index_array[global_index] : global_index;
ccl_gpu_kernel_call(integrator_shade_dedicated_light(NULL, state, render_buffer));
}
}
ccl_gpu_kernel_postfix
ccl_gpu_kernel_threads(GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE)
ccl_gpu_kernel_signature(integrator_queued_paths_array,
int num_states,

16
intern/cycles/kernel/device/hiprt/hiprt_kernels.h
View File

@ -66,6 +66,22 @@ ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
ccl_gpu_kernel_call(integrator_intersect_volume_stack(kg, state));
}
}
ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
ccl_gpu_kernel_signature(integrator_intersect_dedicated_light,
ccl_global const int *path_index_array,
const int work_size,
ccl_global int *stack_buffer)
{
const int global_index = ccl_gpu_global_id_x();
if (global_index < work_size) {
HIPRT_INIT_KERNEL_GLOBAL()
const int state = (path_index_array) ? path_index_array[global_index] : global_index;
ccl_gpu_kernel_call(integrator_intersect_dedicated_light(kg, state));
}
}
ccl_gpu_kernel_postfix
ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
ccl_gpu_kernel_signature(integrator_shade_surface_raytrace,

20
intern/cycles/kernel/device/oneapi/kernel.cpp
View File

@ -122,6 +122,7 @@ size_t oneapi_kernel_preferred_local_size(SyclQueue *queue,
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
@ -129,6 +130,7 @@ size_t oneapi_kernel_preferred_local_size(SyclQueue *queue,
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
preferred_work_group_size = preferred_work_group_size_intersect_shading;
break;
@ -435,6 +437,15 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
oneapi_kernel_integrator_intersect_volume_stack);
break;
}
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT: {
oneapi_call(kg,
cgh,
global_size,
local_size,
args,
oneapi_kernel_integrator_intersect_dedicated_light);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND: {
oneapi_call(
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_background);
@ -474,6 +485,15 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_volume);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT: {
oneapi_call(kg,
cgh,
global_size,
local_size,
args,
oneapi_kernel_integrator_shade_dedicated_light);
break;
}
case DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY: {
oneapi_call(
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_queued_paths_array);

12
intern/cycles/kernel/device/optix/bvh.h
View File

@ -183,6 +183,12 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
return optixIgnoreIntersection();
}
# ifdef __SHADOW_LINKING__
if (intersection_skip_shadow_link(nullptr, ray, object)) {
return optixIgnoreIntersection();
}
# endif
# ifndef __TRANSPARENT_SHADOWS__
/* No transparent shadows support compiled in, make opaque. */
optixSetPayload_5(true);
@ -326,6 +332,12 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
ccl_private Ray *const ray = get_payload_ptr_6<Ray>();
if (visibility & PATH_RAY_SHADOW_OPAQUE) {
#ifdef __SHADOW_LINKING__
if (intersection_skip_shadow_link(nullptr, ray, object)) {
return optixIgnoreIntersection();
}
#endif
if (intersection_skip_self_shadow(ray->self, object, prim)) {
return optixIgnoreIntersection();
}

9
intern/cycles/kernel/device/optix/kernel.cu
View File

@ -18,6 +18,7 @@
#include "kernel/integrator/intersect_shadow.h"
#include "kernel/integrator/intersect_subsurface.h"
#include "kernel/integrator/intersect_volume_stack.h"
#include "kernel/integrator/intersect_dedicated_light.h"
// clang-format on
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
@ -56,3 +57,11 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_st
integrator_intersect_volume_stack(nullptr, path_index);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_dedicated_light()
{
const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ?
kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_dedicated_light(nullptr, path_index);
}

10
intern/cycles/kernel/device/optix/kernel_osl.cu
View File

@ -9,6 +9,7 @@
#include "kernel/bake/bake.h"
#include "kernel/integrator/shade_background.h"
#include "kernel/integrator/shade_dedicated_light.h"
#include "kernel/integrator/shade_light.h"
#include "kernel/integrator/shade_shadow.h"
#include "kernel/integrator/shade_volume.h"
@ -58,6 +59,15 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_shade_shadow()
integrator_shade_shadow(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_dedicated_light()
{
const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ?
kernel_params.path_index_array[global_index] :
global_index;
integrator_shade_dedicated_light(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_shader_eval_displace()
{
KernelShaderEvalInput *const input = (KernelShaderEvalInput *)kernel_params.path_index_array;

22
intern/cycles/kernel/integrator/guiding.h
View File

@ -399,9 +399,8 @@ ccl_device_forceinline void guiding_record_background(KernelGlobals kg,
#endif
}
/* Records the scattered contribution of a next event estimation
* (i.e., a direct light estimate scattered at the current path vertex
* towards the previous vertex). */
/* Records direct lighting from either next event estimation or a dedicated BSDF
* sampled shadow ray. */
ccl_device_forceinline void guiding_record_direct_light(KernelGlobals kg,
IntegratorShadowState state)
{
@ -414,7 +413,22 @@ ccl_device_forceinline void guiding_record_direct_light(KernelGlobals kg,
INTEGRATOR_STATE(state, shadow_path, unlit_throughput));
const float3 Lo_rgb = spectrum_to_rgb(Lo);
openpgl::cpp::AddScatteredContribution(state->shadow_path.path_segment, guiding_vec3f(Lo_rgb));
const float mis_weight = INTEGRATOR_STATE(state, shadow_path, guiding_mis_weight);
if (mis_weight == 0.0f) {
/* Scattered contribution of a next event estimation (i.e., a direct light estimate
* scattered at the current path vertex towards the previous vertex). */
openpgl::cpp::AddScatteredContribution(state->shadow_path.path_segment,
guiding_vec3f(Lo_rgb));
}
else {
/* Dedicated shadow ray for BSDF sampled ray direction.
* The mis weight was already folded into the throughput, so need to divide it out. */
openpgl::cpp::SetDirectContribution(state->shadow_path.path_segment,
guiding_vec3f(Lo_rgb / mis_weight));
openpgl::cpp::SetMiWeight(state->shadow_path.path_segment, mis_weight);
}
}
#endif
}

1
intern/cycles/kernel/integrator/intersect_closest.h
View File

@ -353,6 +353,7 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg,
ray.self.prim = last_isect_prim;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
bool hit = scene_intersect(kg, &ray, visibility, &isect);
/* TODO: remove this and do it in the various intersection functions instead. */

205
intern/cycles/kernel/integrator/intersect_dedicated_light.h Normal file
View File

@ -0,0 +1,205 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2023 Blender Foundation */
#pragma once
#include "kernel/bvh/bvh.h"
#include "kernel/integrator/path_state.h"
#include "kernel/integrator/shade_surface.h"
#include "kernel/integrator/shadow_linking.h"
#include "kernel/light/light.h"
#include "kernel/sample/lcg.h"
CCL_NAMESPACE_BEGIN
#ifdef __SHADOW_LINKING__
# define SHADOW_LINK_MAX_INTERSECTION_COUNT 1024
/* Intersect mesh objects.
*
* Returns the total number of emissive surfaces hit, and the intersection contains a random
* intersected emitter to which the dedicated shadow ray is to eb traced.
*
* NOTE: Sets the ray tmax to the maximum intersection distance (past which no lights are to be
* considered for shadow linking). */
ccl_device int shadow_linking_pick_mesh_intersection(KernelGlobals kg,
IntegratorState state,
ccl_private Ray *ccl_restrict ray,
const int object_receiver,
ccl_private Intersection *ccl_restrict
linked_isect,
ccl_private uint *lcg_state,
int num_hits)
{
/* The tmin will be offset, so store its current value and restore later on, allowing a separate
* light intersection loop starting from the actual ray origin. */
const float old_tmin = ray->tmin;
const uint visibility = path_state_ray_visibility(state);
for (int i = 0; i < SHADOW_LINK_MAX_INTERSECTION_COUNT; i++) {
Intersection current_isect ccl_optional_struct_init;
current_isect.object = OBJECT_NONE;
current_isect.prim = PRIM_NONE;
const bool hit = scene_intersect(kg, ray, visibility, &current_isect);
if (!hit) {
break;
}
/* Only record primitives that potentially have emission.
* TODO: optimize with a dedicated ray visiblity flag, which could then also be
* used once lights are in the BVH as geometry? */
const int shader = intersection_get_shader(kg, &current_isect);
const int shader_flags = kernel_data_fetch(shaders, shader).flags;
if (light_link_object_match(kg, object_receiver, current_isect.object) &&
(shader_flags & SD_HAS_EMISSION))
{
const uint64_t set_membership =
kernel_data_fetch(objects, current_isect.object).shadow_set_membership;
if (set_membership != LIGHT_LINK_MASK_ALL) {
++num_hits;
if ((linked_isect->prim == PRIM_NONE) && (lcg_step_float(lcg_state) < 1.0f / num_hits)) {
*linked_isect = current_isect;
}
}
}
const uint blocker_set = kernel_data_fetch(objects, current_isect.object).blocker_shadow_set;
if (blocker_set == 0) {
/* Contribution from the lights past the default blocker is accumulated using the main path.
*/
ray->tmax = current_isect.t;
break;
}
/* Move the ray forward. */
ray->tmin = intersection_t_offset(current_isect.t);
}
ray->tmin = old_tmin;
return num_hits;
}
/* Pick a light for tracing a shadow ray for the shadow linking.
* Picks a random light which is intersected by the given ray, and stores the intersection result.
* If no lights were hit false is returned.
*
* NOTE: Sets the ray tmax to the maximum intersection distance (past which no lights are to be
* considered for shadow linking). */
ccl_device bool shadow_linking_pick_light_intersection(KernelGlobals kg,
IntegratorState state,
ccl_private Ray *ccl_restrict ray,
ccl_private Intersection *ccl_restrict
linked_isect)
{
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
const int last_type = INTEGRATOR_STATE(state, isect, type);
const int object_receiver = light_link_receiver_forward(kg, state);
uint lcg_state = lcg_state_init(INTEGRATOR_STATE(state, path, rng_hash),
INTEGRATOR_STATE(state, path, rng_offset),
INTEGRATOR_STATE(state, path, sample),
0x68bc21eb);
/* Indicate that no intersection has been picked yet. */
linked_isect->prim = PRIM_NONE;
int num_hits = 0;
// TODO: Only if there are emissive meshes in the scene?
// TODO: Only if the ray hits any light? As in, check that there is a light first, before
// tracing potentially expensive ray.
num_hits = shadow_linking_pick_mesh_intersection(
kg, state, ray, object_receiver, linked_isect, &lcg_state, num_hits);
num_hits = lights_intersect_shadow_linked(kg,
ray,
linked_isect,
ray->self.prim,
ray->self.object,
last_type,
path_flag,
object_receiver,
&lcg_state,
num_hits);
if (num_hits == 0) {
return false;
}
INTEGRATOR_STATE_WRITE(state, shadow_link, dedicated_light_weight) = num_hits;
return true;
}
/* Check whether a special shadow ray is needed to calculate direct light contribution which comes
* from emitters which are behind objects which are blocking light for the main path, but are
* excluded from blocking light via shadow linking.
*
* If a special ray is needed a blocked light kernel is scheduled and true is returned, otherwise
* false is returned. */
ccl_device bool shadow_linking_intersect(KernelGlobals kg, IntegratorState state)
{
/* Verify that the kernel is only scheduled if it is actually needed. */
kernel_assert(shadow_linking_scene_need_shadow_ray(kg, state));
/* Read ray from integrator state into local memory. */
Ray ray ccl_optional_struct_init;
integrator_state_read_ray(state, &ray);
ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
ray.self.object = INTEGRATOR_STATE(state, isect, object);
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
Intersection isect ccl_optional_struct_init;
if (!shadow_linking_pick_light_intersection(kg, state, &ray, &isect)) {
/* No light is hit, no need in the extra shadow ray for the direct light. */
return false;
}
/* Make a copy of primitives needed by the main path self-intersection check before writing the
* new intersection. Those primitives will be restored before the main path is returned to the
* intersect_closest state. */
shadow_linking_store_last_primitives(state);
/* Write intersection result into global integrator state memory, so that the
* shade_dedicated_light kernel can use it for calculation of the light sample. */
integrator_state_write_isect(state, &isect);
integrator_path_next(kg,
state,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT,
DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT);
return true;
}
#endif /* __SHADOW_LINKING__ */
ccl_device void integrator_intersect_dedicated_light(KernelGlobals kg, IntegratorState state)
{
PROFILING_INIT(kg, PROFILING_INTERSECT_DEDICATED_LIGHT);
#ifdef __SHADOW_LINKING__
if (shadow_linking_intersect(kg, state)) {
return;
}
#else
kernel_assert(!"integrator_intersect_dedicated_light is not supposed to be scheduled");
#endif
integrator_shade_surface_next_kernel<DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT>(kg,
state);
}
CCL_NAMESPACE_END

8
intern/cycles/kernel/integrator/intersect_shadow.h
View File

@ -34,6 +34,9 @@ ccl_device bool integrate_intersect_shadow_opaque(KernelGlobals kg,
Intersection isect;
const bool opaque_hit = scene_intersect(kg, ray, visibility & opaque_mask, &isect);
/* Only record the number of hits if nothing was hit, so that the shadow shading kernel does not
* consider any intersections. There is no need to write anything to the state if the hit is
* opaque because in this case the path is terminated. */
if (!opaque_hit) {
INTEGRATOR_STATE_WRITE(state, shadow_path, num_hits) = 0;
}
@ -144,10 +147,7 @@ ccl_device void integrator_intersect_shadow(KernelGlobals kg, IntegratorShadowSt
/* Read ray from integrator state into local memory. */
Ray ray ccl_optional_struct_init;
integrator_state_read_shadow_ray(state, &ray);
ray.self.object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, object);
ray.self.prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, prim);
ray.self.light_object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, object);
ray.self.light_prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, prim);
integrator_state_read_shadow_ray_self(kg, state, &ray);
/* Compute visibility. */
const uint visibility = integrate_intersect_shadow_visibility(kg, state);

2
intern/cycles/kernel/integrator/intersect_volume_stack.h
View File

@ -29,6 +29,7 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
volume_ray.self.light_object = OBJECT_NONE;
volume_ray.self.light_prim = PRIM_NONE;
volume_ray.self.light = LAMP_NONE;
/* Store to avoid global fetches on every intersection step. */
const uint volume_stack_size = kernel_data.volume_stack_size;
@ -84,6 +85,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
volume_ray.self.prim = PRIM_NONE;
volume_ray.self.light_object = OBJECT_NONE;
volume_ray.self.light_prim = PRIM_NONE;
volume_ray.self.light = LAMP_NONE;
int stack_index = 0, enclosed_index = 0;

8
intern/cycles/kernel/integrator/megakernel.h
View File

@ -5,10 +5,12 @@
#include "kernel/integrator/init_from_camera.h"
#include "kernel/integrator/intersect_closest.h"
#include "kernel/integrator/intersect_dedicated_light.h"
#include "kernel/integrator/intersect_shadow.h"
#include "kernel/integrator/intersect_subsurface.h"
#include "kernel/integrator/intersect_volume_stack.h"
#include "kernel/integrator/shade_background.h"
#include "kernel/integrator/shade_dedicated_light.h"
#include "kernel/integrator/shade_light.h"
#include "kernel/integrator/shade_shadow.h"
#include "kernel/integrator/shade_surface.h"
@ -83,12 +85,18 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
integrator_shade_light(kg, state, render_buffer);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
integrator_shade_dedicated_light(kg, state, render_buffer);
break;
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
integrator_intersect_subsurface(kg, state);
break;
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
integrator_intersect_volume_stack(kg, state);
break;
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
integrator_intersect_dedicated_light(kg, state);
break;
default:
kernel_assert(0);
break;

4
intern/cycles/kernel/integrator/mnee.h
View File

@ -408,6 +408,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg,
Ray projection_ray;
projection_ray.self.light_object = OBJECT_NONE;
projection_ray.self.light_prim = PRIM_NONE;
projection_ray.self.light = LAMP_NONE;
projection_ray.dP = differential_make_compact(sd->dP);
projection_ray.dD = differential_zero_compact();
projection_ray.tmin = 0.0f;
@ -482,6 +483,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg,
if (!hit)
break;
// TODO: Is the fetch needed here? The shade_surface simply reads isect->object.
int hit_object = (projection_isect.object == OBJECT_NONE) ?
kernel_data_fetch(prim_object, projection_isect.prim) :
projection_isect.object;
@ -826,6 +828,7 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg,
Ray probe_ray;
probe_ray.self.light_object = ls->object;
probe_ray.self.light_prim = ls->prim;
probe_ray.self.light = ls->lamp;
probe_ray.tmin = 0.0f;
probe_ray.dP = differential_make_compact(sd->dP);
probe_ray.dD = differential_zero_compact();
@ -926,6 +929,7 @@ ccl_device_forceinline int kernel_path_mnee_sample(KernelGlobals kg,
probe_ray.self.prim = sd->prim;
probe_ray.self.light_object = ls->object;
probe_ray.self.light_prim = ls->prim;
probe_ray.self.light = ls->lamp;
probe_ray.P = sd->P;
probe_ray.tmin = 0.0f;
if (ls->t == FLT_MAX) {

6
intern/cycles/kernel/integrator/path_state.h
View File

@ -89,6 +89,12 @@ ccl_device_inline void path_state_init_integrator(KernelGlobals kg,
INTEGRATOR_STATE_WRITE(state, path, denoising_feature_throughput) = one_spectrum();
}
#endif
#ifdef __LIGHT_LINKING__
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING) {
INTEGRATOR_STATE_WRITE(state, path, mis_ray_object) = OBJECT_NONE;
}
#endif
}
ccl_device_inline void path_state_next(KernelGlobals kg,

31
intern/cycles/kernel/integrator/shade_background.h
View File

@ -22,14 +22,8 @@ ccl_device Spectrum integrator_eval_background_shader(KernelGlobals kg,
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
/* Use visibility flag to skip lights. */
if (shader & SHADER_EXCLUDE_ANY) {
if (((shader & SHADER_EXCLUDE_DIFFUSE) && (path_flag & PATH_RAY_DIFFUSE)) ||
((shader & SHADER_EXCLUDE_GLOSSY) && ((path_flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
(PATH_RAY_GLOSSY | PATH_RAY_REFLECT))) ||
((shader & SHADER_EXCLUDE_TRANSMIT) && (path_flag & PATH_RAY_TRANSMIT)) ||
((shader & SHADER_EXCLUDE_CAMERA) && (path_flag & PATH_RAY_CAMERA)) ||
((shader & SHADER_EXCLUDE_SCATTER) && (path_flag & PATH_RAY_VOLUME_SCATTER)))
return zero_spectrum();
if (!is_light_shader_visible_to_path(shader, path_flag)) {
return zero_spectrum();
}
/* Use fast constant background color if available. */
@ -140,16 +134,19 @@ ccl_device_inline void integrate_distant_lights(KernelGlobals kg,
/* Use visibility flag to skip lights. */
#ifdef __PASSES__
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
continue;
}
#endif
if (ls.shader & SHADER_EXCLUDE_ANY) {
if (((ls.shader & SHADER_EXCLUDE_DIFFUSE) && (path_flag & PATH_RAY_DIFFUSE)) ||
((ls.shader & SHADER_EXCLUDE_GLOSSY) &&
((path_flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
(PATH_RAY_GLOSSY | PATH_RAY_REFLECT))) ||
((ls.shader & SHADER_EXCLUDE_TRANSMIT) && (path_flag & PATH_RAY_TRANSMIT)) ||
((ls.shader & SHADER_EXCLUDE_CAMERA) && (path_flag & PATH_RAY_CAMERA)) ||
((ls.shader & SHADER_EXCLUDE_SCATTER) && (path_flag & PATH_RAY_VOLUME_SCATTER)))
continue;
#ifdef __LIGHT_LINKING__
if (!light_link_light_match(kg, light_link_receiver_forward(kg, state), lamp)) {
continue;
}
#endif
#ifdef __SHADOW_LINKING__
if (kernel_data_fetch(lights, lamp).shadow_set_membership != LIGHT_LINK_MASK_ALL) {
continue;
}
#endif

271
intern/cycles/kernel/integrator/shade_dedicated_light.h Normal file
View File

@ -0,0 +1,271 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2023 Blender Foundation */
#pragma once
#include "kernel/integrator/path_state.h"
#include "kernel/integrator/shade_surface.h"
#include "kernel/light/distant.h"
#include "kernel/light/light.h"
CCL_NAMESPACE_BEGIN
#ifdef __SHADOW_LINKING__
ccl_device_inline bool shadow_linking_light_sample_from_intersection(
KernelGlobals kg,
ccl_private const Intersection &ccl_restrict isect,
ccl_private const Ray &ccl_restrict ray,
ccl_private LightSample *ccl_restrict ls)
{
const int lamp = isect.prim;
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
const LightType type = LightType(klight->type);
if (type == LIGHT_DISTANT) {
return distant_light_sample_from_intersection(kg, ray.D, lamp, ls);
}
return light_sample_from_intersection(kg, &isect, ray.P, ray.D, ls);
}
ccl_device_inline float shadow_linking_light_sample_mis_weight(KernelGlobals kg,
IntegratorState state,
const uint32_t path_flag,
const ccl_private LightSample *ls,
const float3 P)
{
if (ls->type == LIGHT_DISTANT) {
return light_sample_mis_weight_forward_distant(kg, state, path_flag, ls);
}
return light_sample_mis_weight_forward_lamp(kg, state, path_flag, ls, P);
}
/* Setup ray for the shadow path.
* Expects that the current state of the ray is the one calculated by the surface bounce, and the
* intersection corresponds to a point on an emitter. */
ccl_device void shadow_linking_setup_ray_from_intersection(
IntegratorState state,
ccl_private Ray *ccl_restrict ray,
ccl_private const Intersection *ccl_restrict isect)
{
/* The ray->tmin follows the value configured at the surface bounce.
* it is the same for the continued main path and for this shadow ray. There is no need to push
* it forward here. */
ray->tmax = isect->t;
/* Use the same self intersection primitives as the main path.
* Those are copied to the dedicated storage from the main intersection after the surface bounce,
* but before the main intersection is re-used to find light to trace a ray to. */
ray->self.object = INTEGRATOR_STATE(state, shadow_link, last_isect_object);
ray->self.prim = INTEGRATOR_STATE(state, shadow_link, last_isect_prim);
if (isect->type == PRIMITIVE_LAMP) {
ray->self.light_object = OBJECT_NONE;
ray->self.light_prim = PRIM_NONE;
ray->self.light = isect->prim;
}
else {
ray->self.light_object = isect->object;
ray->self.light_prim = isect->prim;
ray->self.light = LAMP_NONE;
}
}
ccl_device bool shadow_linking_shade_light(KernelGlobals kg,
IntegratorState state,
ccl_private Ray &ccl_restrict ray,
ccl_private Intersection &ccl_restrict isect,
ccl_private ShaderData *emission_sd,
ccl_private Spectrum &ccl_restrict bsdf_spectrum,
ccl_private float &mis_weight,
ccl_private int &ccl_restrict light_group)
{
LightSample ls ccl_optional_struct_init;
const bool use_light_sample = shadow_linking_light_sample_from_intersection(kg, isect, ray, &ls);
if (!use_light_sample) {
/* No light to be sampled, so no direct light contribution either. */
return false;
}
const Spectrum light_eval = light_sample_shader_eval(kg, state, emission_sd, &ls, ray.time);
if (is_zero(light_eval)) {
return false;
}
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
return false;
}
/* MIS weighting. */
if (!(path_flag & PATH_RAY_MIS_SKIP)) {
mis_weight = shadow_linking_light_sample_mis_weight(kg, state, path_flag, &ls, ray.P);
}
bsdf_spectrum = light_eval * mis_weight *
INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
// TODO(: De-duplicate with the shade_surface.
// Possibly by ensuring ls->group is always assigned properly.
light_group = ls.type != LIGHT_BACKGROUND ? ls.group : kernel_data.background.lightgroup;
return true;
}
ccl_device bool shadow_linking_shade_surface_emission(KernelGlobals kg,
IntegratorState state,
ccl_private Ray &ccl_restrict ray,
ccl_private Intersection &ccl_restrict isect,
ccl_private ShaderData *emission_sd,
ccl_global float *ccl_restrict render_buffer,
ccl_private Spectrum &ccl_restrict
bsdf_spectrum,
ccl_private float &mis_weight,
ccl_private int &ccl_restrict light_group)
{
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
integrate_surface_shader_setup(kg, state, emission_sd);
# ifdef __VOLUME__
if (emission_sd->flag & SD_HAS_ONLY_VOLUME) {
return false;
}
# endif
surface_shader_eval<KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT>(
kg, state, emission_sd, render_buffer, path_flag);
surface_shader_prepare_closures(kg, state, emission_sd, path_flag);
if ((emission_sd->flag & SD_EMISSION) == 0) {
return false;
}
const Spectrum L = surface_shader_emission(emission_sd);
const bool has_mis = !(path_flag & PATH_RAY_MIS_SKIP) &&
(emission_sd->flag &
((emission_sd->flag & SD_BACKFACING) ? SD_MIS_BACK : SD_MIS_FRONT));
# ifdef __HAIR__
if (has_mis && (emission_sd->type & PRIMITIVE_TRIANGLE))
# else
if (has_mis)
# endif
{
mis_weight = light_sample_mis_weight_forward_surface(kg, state, path_flag, emission_sd);
}
bsdf_spectrum = L * mis_weight * INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
light_group = object_lightgroup(kg, emission_sd->object);
return true;
}
ccl_device void shadow_linking_shade(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
{
/* Read intersection from integrator state into local memory. */
Intersection isect ccl_optional_struct_init;
integrator_state_read_isect(state, &isect);
/* Read ray from integrator state into local memory. */
Ray ray ccl_optional_struct_init;
integrator_state_read_ray(state, &ray);
ShaderDataCausticsStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
Spectrum bsdf_spectrum;
float mis_weight = 1.0f;
int light_group = LIGHTGROUP_NONE;
if (isect.type == PRIMITIVE_LAMP) {
if (!shadow_linking_shade_light(
kg, state, ray, isect, emission_sd, bsdf_spectrum, mis_weight, light_group))
{
return;
}
}
else {
if (!shadow_linking_shade_surface_emission(kg,
state,
ray,
isect,
emission_sd,
render_buffer,
bsdf_spectrum,
mis_weight,
light_group))
{
return;
}
}
if (is_zero(bsdf_spectrum)) {
return;
}
shadow_linking_setup_ray_from_intersection(state, &ray, &isect);
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
kg, state, &ray, bsdf_spectrum, 0, light_group);
/* The light is accumulated from the shade_surface kernel, which will make the clamping decision
* based on the actual value of the bounce. For the dedicated shadow ray we want to follow the
* main path clamping rules, which subtracts one from the bounds before accumulation. */
INTEGRATOR_STATE_WRITE(
shadow_state, shadow_path, bounce) = INTEGRATOR_STATE(shadow_state, shadow_path, bounce) - 1;
/* No need to update the volume stack as the surface bounce already performed enter-exit check.
*/
const uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
/* The diffuse and glossy pass weights are written into the main path as part of the path
* configuration at a surface bounce. */
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = INTEGRATOR_STATE(
state, path, pass_diffuse_weight);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = INTEGRATOR_STATE(
state, path, pass_glossy_weight);
}
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, flag) = shadow_flag;
# ifdef __PATH_GUIDING__
if (kernel_data.integrator.train_guiding) {
guiding_record_light_surface_segment(kg, state, &isect);
INTEGRATOR_STATE(shadow_state, shadow_path, guiding_mis_weight) = mis_weight;
}
# endif
}
#endif /* __SHADOW_LINKING__ */
ccl_device void integrator_shade_dedicated_light(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
{
PROFILING_INIT(kg, PROFILING_SHADE_DEDICATED_LIGHT);
#ifdef __SHADOW_LINKING__
shadow_linking_shade(kg, state, render_buffer);
/* Restore self-intersection check primitives in the main state before returning to the
* intersect_closest() state. */
shadow_linking_restore_last_primitives(state);
#else
kernel_assert(!"integrator_intersect_dedicated_light is not supposed to be scheduled");
#endif
integrator_shade_surface_next_kernel<DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT>(kg, state);
}
CCL_NAMESPACE_END

11
intern/cycles/kernel/integrator/shade_light.h
View File

@ -37,15 +37,8 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
/* Use visibility flag to skip lights. */
#ifdef __PASSES__
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
if (ls.shader & SHADER_EXCLUDE_ANY) {
if (((ls.shader & SHADER_EXCLUDE_DIFFUSE) && (path_flag & PATH_RAY_DIFFUSE)) ||
((ls.shader & SHADER_EXCLUDE_GLOSSY) &&
((path_flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
(PATH_RAY_GLOSSY | PATH_RAY_REFLECT))) ||
((ls.shader & SHADER_EXCLUDE_TRANSMIT) && (path_flag & PATH_RAY_TRANSMIT)) ||
((ls.shader & SHADER_EXCLUDE_SCATTER) && (path_flag & PATH_RAY_VOLUME_SCATTER)))
return;
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
return;
}
#endif

1
intern/cycles/kernel/integrator/shade_shadow.h
View File

@ -75,6 +75,7 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg,
ray.self.prim = PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
/* Modify ray position and length to match current segment. */
ray.tmin = (hit == 0) ? ray.tmin : INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t);
ray.tmax = (hit < num_recorded_hits) ? INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit, t) :

246
intern/cycles/kernel/integrator/shade_surface.h
View File

@ -3,6 +3,9 @@
#pragma once
#include "kernel/integrator/path_state.h"
#include "kernel/integrator/surface_shader.h"
#include "kernel/film/data_passes.h"
#include "kernel/film/denoising_passes.h"
#include "kernel/film/light_passes.h"
@ -10,9 +13,8 @@
#include "kernel/integrator/mnee.h"
#include "kernel/integrator/guiding.h"
#include "kernel/integrator/path_state.h"
#include "kernel/integrator/shadow_linking.h"
#include "kernel/integrator/subsurface.h"
#include "kernel/integrator/surface_shader.h"
#include "kernel/integrator/volume_stack.h"
#include "kernel/light/sample.h"
@ -111,8 +113,26 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
ccl_global float *ccl_restrict
render_buffer)
{
#ifdef __LIGHT_LINKING__
if (!light_link_object_match(kg, light_link_receiver_forward(kg, state), sd->object)) {
return;
}
#endif
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
#ifdef __SHADOW_LINKING__
/* Indirect emission of shadow-linked emissive surfaces is done via shadow rays to dedicated
* light sources. */
if (kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING) {
if (!(path_flag & PATH_RAY_CAMERA) &&
kernel_data_fetch(objects, sd->object).shadow_set_membership != LIGHT_LINK_MASK_ALL)
{
return;
}
}
#endif
/* Evaluate emissive closure. */
Spectrum L = surface_shader_emission(sd);
float mis_weight = 1.0f;
@ -134,6 +154,84 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
kg, state, L, mis_weight, render_buffer, object_lightgroup(kg, sd->object));
}
/* Branch off a shadow path and initialize common part of it.
* THe common is between the surface shading and configuration of a special shadow ray for the
* shadow linking. */
ccl_device_inline IntegratorShadowState
integrate_direct_light_shadow_init_common(KernelGlobals kg,
IntegratorState state,
ccl_private const Ray *ccl_restrict ray,
const Spectrum bsdf_spectrum,
const int light_group,
const int mnee_vertex_count)
{
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
/* Copy volume stack and enter/exit volume. */
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(shadow_state, ray);
integrator_state_write_shadow_ray_self(kg, shadow_state, ray);
/* Copy state from main path to shadow path. */
const Spectrum unlit_throughput = INTEGRATOR_STATE(state, path, throughput);
const Spectrum throughput = unlit_throughput * bsdf_spectrum;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
state, path, render_pixel_index);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(
state, path, rng_offset);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_hash) = INTEGRATOR_STATE(
state, path, rng_hash);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, sample) = INTEGRATOR_STATE(
state, path, sample);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transparent_bounce) = INTEGRATOR_STATE(
state, path, transparent_bounce);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, glossy_bounce) = INTEGRATOR_STATE(
state, path, glossy_bounce);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput;
#ifdef __MNEE__
if (mnee_vertex_count > 0) {
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) =
INTEGRATOR_STATE(state, path, transmission_bounce) + mnee_vertex_count - 1;
INTEGRATOR_STATE_WRITE(shadow_state,
shadow_path,
diffuse_bounce) = INTEGRATOR_STATE(state, path, diffuse_bounce) + 1;
INTEGRATOR_STATE_WRITE(shadow_state,
shadow_path,
bounce) = INTEGRATOR_STATE(state, path, bounce) + mnee_vertex_count;
}
else
#endif
{
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) = INTEGRATOR_STATE(
state, path, transmission_bounce);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, diffuse_bounce) = INTEGRATOR_STATE(
state, path, diffuse_bounce);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bounce) = INTEGRATOR_STATE(
state, path, bounce);
}
/* Write Lightgroup, +1 as lightgroup is int but we need to encode into a uint8_t. */
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, lightgroup) = light_group;
#ifdef __PATH_GUIDING__
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, unlit_throughput) = unlit_throughput;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, path_segment) = INTEGRATOR_STATE(
state, guiding, path_segment);
INTEGRATOR_STATE(shadow_state, shadow_path, guiding_mis_weight) = 0.0f;
#endif
return shadow_state;
}
/* Path tracing: sample point on light and evaluate light shader, then
* queue shadow ray to be traced. */
template<uint node_feature_mask>
@ -162,6 +260,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
sd->time,
sd->P,
sd->N,
light_link_receiver_nee(kg, sd),
sd->flag,
bounce,
path_flag,
@ -241,12 +340,18 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
light_sample_to_surface_shadow_ray(kg, sd, &ls, &ray);
}
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
if (ray.self.object != OBJECT_NONE) {
ray.P = integrate_surface_ray_offset(kg, sd, ray.P, ray.D);
}
/* Copy volume stack and enter/exit volume. */
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
/* Branch off shadow kernel. */
// TODO(: De-duplicate with the shade_Dedicated_light.
// Possibly by ensuring ls->group is always assigned properly.
const int light_group = ls.type != LIGHT_BACKGROUND ? ls.group :
kernel_data.background.lightgroup;
IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
kg, state, &ray, bsdf_eval_sum(&bsdf_eval), mnee_vertex_count, light_group);
if (is_transmission) {
#ifdef __VOLUME__
@ -254,22 +359,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
#endif
}
if (ray.self.object != OBJECT_NONE) {
ray.P = integrate_surface_ray_offset(kg, sd, ray.P, ray.D);
}
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(shadow_state, &ray);
// Save memory by storing the light and object indices in the shadow_isect
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
/* Copy state from main path to shadow path. */
uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
const Spectrum unlit_throughput = INTEGRATOR_STATE(state, path, throughput);
const Spectrum throughput = unlit_throughput * bsdf_eval_sum(&bsdf_eval);
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
PackedSpectrum pass_diffuse_weight;
@ -291,55 +381,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
}
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, render_pixel_index) = INTEGRATOR_STATE(
state, path, render_pixel_index);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_offset) = INTEGRATOR_STATE(
state, path, rng_offset);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, rng_hash) = INTEGRATOR_STATE(
state, path, rng_hash);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, sample) = INTEGRATOR_STATE(
state, path, sample);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, flag) = shadow_flag;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transparent_bounce) = INTEGRATOR_STATE(
state, path, transparent_bounce);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, glossy_bounce) = INTEGRATOR_STATE(
state, path, glossy_bounce);
#ifdef __MNEE__
if (mnee_vertex_count > 0) {
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) =
INTEGRATOR_STATE(state, path, transmission_bounce) + mnee_vertex_count - 1;
INTEGRATOR_STATE_WRITE(shadow_state,
shadow_path,
diffuse_bounce) = INTEGRATOR_STATE(state, path, diffuse_bounce) + 1;
INTEGRATOR_STATE_WRITE(shadow_state,
shadow_path,
bounce) = INTEGRATOR_STATE(state, path, bounce) + mnee_vertex_count;
}
else
#endif
{
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, transmission_bounce) = INTEGRATOR_STATE(
state, path, transmission_bounce);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, diffuse_bounce) = INTEGRATOR_STATE(
state, path, diffuse_bounce);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bounce) = INTEGRATOR_STATE(
state, path, bounce);
}
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = throughput;
/* Write Lightgroup, +1 as lightgroup is int but we need to encode into a uint8_t. */
INTEGRATOR_STATE_WRITE(
shadow_state, shadow_path, lightgroup) = (ls.type != LIGHT_BACKGROUND) ?
ls.group + 1 :
kernel_data.background.lightgroup + 1;
#ifdef __PATH_GUIDING__
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, unlit_throughput) = unlit_throughput;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, path_segment) = INTEGRATOR_STATE(
state, guiding, path_segment);
#endif
}
/* Path tracing: bounce off or through surface with new direction. */
@ -453,6 +495,11 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
unguided_bsdf_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf));
}
#ifdef __LIGHT_LINKING__
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING) {
INTEGRATOR_STATE_WRITE(state, path, mis_ray_object) = sd->object;
}
#endif
path_state_next(kg, state, label, sd->flag);
@ -546,6 +593,7 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
ray.self.prim = (skip_self) ? sd->prim : PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
ray.dP = differential_zero_compact();
ray.dD = differential_zero_compact();
@ -558,10 +606,7 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(shadow_state, &ray);
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
integrator_state_write_shadow_ray_self(kg, shadow_state, &ray);
/* Copy state from main path to shadow path. */
const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
@ -590,9 +635,9 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
#endif /* defined(__AO__) */
template<uint node_feature_mask>
ccl_device bool integrate_surface(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
ccl_device int integrate_surface(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
{
PROFILING_INIT_FOR_SHADER(kg, PROFILING_SHADE_SURFACE_SETUP);
@ -645,7 +690,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
/* Evaluate holdout. */
if (!integrate_surface_holdout(kg, state, &sd, render_buffer)) {
return false;
return LABEL_NONE;
}
/* Write emission. */
@ -659,7 +704,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
*
* Also ensure we don't do it twice for SSS at both the entry and exit point. */
if (integrate_surface_terminate(state, path_flag)) {
return false;
return LABEL_NONE;
}
/* Write render passes. */
@ -699,7 +744,7 @@ ccl_device bool integrate_surface(KernelGlobals kg,
}
else {
if (integrate_surface_terminate(state, path_flag)) {
return false;
return LABEL_NONE;
}
PROFILING_EVENT(PROFILING_SHADE_SURFACE_INDIRECT_LIGHT);
@ -712,7 +757,20 @@ ccl_device bool integrate_surface(KernelGlobals kg,
}
#endif
return continue_path_label != 0;
return continue_path_label;
}
template<DeviceKernel current_kernel>
ccl_device_forceinline void integrator_shade_surface_next_kernel(KernelGlobals kg,
IntegratorState state)
{
if (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_SUBSURFACE) {
integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE);
}
else {
kernel_assert(INTEGRATOR_STATE(state, ray, tmax) != 0.0f);
integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
}
}
template<uint node_feature_mask = KERNEL_FEATURE_NODE_MASK_SURFACE & ~KERNEL_FEATURE_NODE_RAYTRACE,
@ -721,19 +779,23 @@ ccl_device_forceinline void integrator_shade_surface(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
{
if (integrate_surface<node_feature_mask>(kg, state, render_buffer)) {
if (INTEGRATOR_STATE(state, path, flag) & PATH_RAY_SUBSURFACE) {
integrator_path_next(
kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE);
}
else {
kernel_assert(INTEGRATOR_STATE(state, ray, tmax) != 0.0f);
integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
}
}
else {
const int continue_path_label = integrate_surface<node_feature_mask>(kg, state, render_buffer);
if (continue_path_label == LABEL_NONE) {
integrator_path_terminate(kg, state, current_kernel);
return;
}
#ifdef __SHADOW_LINKING__
/* No need to cast shadow linking rays at a transparent bounce: the lights will be accumulated
* via the main path in this case. */
if ((continue_path_label & LABEL_TRANSPARENT) == 0) {
if (shadow_linking_schedule_intersection_kernel<current_kernel>(kg, state)) {
return;
}
}
#endif
integrator_shade_surface_next_kernel<current_kernel>(kg, state);
}
ccl_device_forceinline void integrator_shade_surface_raytrace(

50
intern/cycles/kernel/integrator/shade_volume.h
View File

@ -10,6 +10,7 @@
#include "kernel/integrator/guiding.h"
#include "kernel/integrator/intersect_closest.h"
#include "kernel/integrator/path_state.h"
#include "kernel/integrator/shadow_linking.h"
#include "kernel/integrator/volume_shader.h"
#include "kernel/integrator/volume_stack.h"
@ -673,8 +674,10 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
/* Write accumulated emission. */
if (!is_zero(accum_emission)) {
film_write_volume_emission(
kg, state, accum_emission, render_buffer, object_lightgroup(kg, sd->object));
if (light_link_object_match(kg, light_link_receiver_forward(kg, state), sd->object)) {
film_write_volume_emission(
kg, state, accum_emission, render_buffer, object_lightgroup(kg, sd->object));
}
}
# ifdef __DENOISING_FEATURES__
@ -714,6 +717,7 @@ ccl_device_forceinline bool integrate_volume_equiangular_sample_light(
sd->P,
ray->D,
ray->tmax - ray->tmin,
light_link_receiver_nee(kg, sd),
bounce,
path_flag,
&ls))
@ -778,6 +782,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
sd->time,
P,
zero_float3(),
light_link_receiver_nee(kg, sd),
SD_BSDF_HAS_TRANSMISSION,
bounce,
path_flag,
@ -831,10 +836,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(shadow_state, &ray);
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ray.self.object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ray.self.prim;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray.self.light_object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = ray.self.light_prim;
integrator_state_write_shadow_ray_self(kg, shadow_state, &ray);
/* Copy state from main path to shadow path. */
const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
@ -891,6 +893,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, unlit_throughput) = unlit_throughput;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, path_segment) = INTEGRATOR_STATE(
state, guiding, path_segment);
INTEGRATOR_STATE(shadow_state, shadow_path, guiding_mis_weight) = 0.0f;
# endif
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
@ -985,6 +988,12 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
unguided_phase_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf));
# ifdef __LIGHT_LINKING__
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING) {
INTEGRATOR_STATE_WRITE(state, path, mis_ray_object) = sd->object;
}
# endif
path_state_next(kg, state, label, sd->flag);
return true;
}
@ -1188,27 +1197,32 @@ ccl_device void integrator_shade_volume(KernelGlobals kg,
volume_stack_clean(kg, state);
}
VolumeIntegrateEvent event = volume_integrate(kg, state, &ray, render_buffer);
if (event == VOLUME_PATH_SCATTERED) {
/* Queue intersect_closest kernel. */
integrator_path_next(kg,
state,
DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
return;
}
else if (event == VOLUME_PATH_MISSED) {
const VolumeIntegrateEvent event = volume_integrate(kg, state, &ray, render_buffer);
if (event == VOLUME_PATH_MISSED) {
/* End path. */
integrator_path_terminate(kg, state, DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME);
return;
}
else {
if (event == VOLUME_PATH_ATTENUATED) {
/* Continue to background, light or surface. */
integrator_intersect_next_kernel_after_volume<DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME>(
kg, state, &isect, render_buffer);
return;
}
# ifdef __SHADOW_LINKING__
if (shadow_linking_schedule_intersection_kernel<DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME>(kg,
state)) {
return;
}
# endif /* __SHADOW_LINKING__ */
/* Queue intersect_closest kernel. */
integrator_path_next(kg,
state,
DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
#endif /* __VOLUME__ */
}

71
intern/cycles/kernel/integrator/shadow_linking.h Normal file
View File

@ -0,0 +1,71 @@
/* SPDX-License-Identifier: Apache-2.0
* Copyright 2011-2023 Blender Foundation */
#pragma once
#include "kernel/integrator/path_state.h"
#include "kernel/integrator/state_util.h"
CCL_NAMESPACE_BEGIN
#ifdef __SHADOW_LINKING__
/* Check whether special shadow rays for shadow linking are needed in the current scene
* configuration. */
ccl_device_forceinline bool shadow_linking_scene_need_shadow_ray(KernelGlobals kg,
IntegratorState state)
{
if (!(kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING)) {
/* No shadow linking in the scene, so no need to trace any extra rays. */
return false;
}
/* The distant lights might be using shadow linking, and they are not counted as
* kernel_data.integrator.use_light_mis.
* So there is a potential to avoid extra rays from being traced, but it requires more granular
* flags set in the integrator. */
return true;
}
/* Shadow linking re-used the main path intersection to store information about the light to which
* the extra ray is to be traced (this intersection communicates light between the shadow blocker
* intersection and shading kernels).
* These utilities makes a copy of the fields from the main intersection which are needed by the
* intersect_closest kernel after the surface bounce. */
ccl_device_forceinline void shadow_linking_store_last_primitives(IntegratorState state)
{
INTEGRATOR_STATE_WRITE(state, shadow_link, last_isect_prim) = INTEGRATOR_STATE(
state, isect, prim);
INTEGRATOR_STATE_WRITE(state, shadow_link, last_isect_object) = INTEGRATOR_STATE(
state, isect, object);
}
ccl_device_forceinline void shadow_linking_restore_last_primitives(IntegratorState state)
{
INTEGRATOR_STATE_WRITE(state, isect, prim) = INTEGRATOR_STATE(
state, shadow_link, last_isect_prim);
INTEGRATOR_STATE_WRITE(state, isect, object) = INTEGRATOR_STATE(
state, shadow_link, last_isect_object);
}
/* Schedule shadow linking intersection kernel if it is needed.
* Returns true if the shadow linking specific kernel has been scheduled, false otherwise. */
template<DeviceKernel current_kernel>
ccl_device_inline bool shadow_linking_schedule_intersection_kernel(KernelGlobals kg,
IntegratorState state)
{
if (!shadow_linking_scene_need_shadow_ray(kg, state)) {
return false;
}
integrator_path_next(
kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT);
return true;
}
#endif /* __SHADOW_LINKING__ */
CCL_NAMESPACE_END

3
intern/cycles/kernel/integrator/shadow_state_template.h
View File

@ -50,6 +50,7 @@ KERNEL_STRUCT_MEMBER(shadow_path,
#else
KERNEL_STRUCT_MEMBER(shadow_path, uint64_t, path_segment, KERNEL_FEATURE_PATH_GUIDING)
#endif
KERNEL_STRUCT_MEMBER(shadow_path, float, guiding_mis_weight, KERNEL_FEATURE_PATH_GUIDING)
KERNEL_STRUCT_END(shadow_path)
/********************************** Shadow Ray *******************************/
@ -61,7 +62,7 @@ KERNEL_STRUCT_MEMBER(shadow_ray, float, tmin, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, tmax, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, time, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, dP, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, int, object, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, int, self_light, KERNEL_FEATURE_SHADOW_LINKING)
KERNEL_STRUCT_END(shadow_ray)
/*********************** Shadow Intersection result **************************/

12
intern/cycles/kernel/integrator/state_template.h
View File

@ -41,6 +41,9 @@ KERNEL_STRUCT_MEMBER(path, uint8_t, mnee, KERNEL_FEATURE_PATH_TRACING)
* zero and distance. Note that transparency and volume attenuation increase
* the ray tmin but keep P unmodified so that this works. */
KERNEL_STRUCT_MEMBER(path, float, mis_ray_pdf, KERNEL_FEATURE_PATH_TRACING)
/* Object at last scatter point for light linking. */
KERNEL_STRUCT_MEMBER(path, int, mis_ray_object, KERNEL_FEATURE_LIGHT_LINKING)
/* Normal at last scatter point for light tree. */
KERNEL_STRUCT_MEMBER(path, packed_float3, mis_origin_n, KERNEL_FEATURE_PATH_TRACING)
/* Filter glossy. */
KERNEL_STRUCT_MEMBER(path, float, min_ray_pdf, KERNEL_FEATURE_PATH_TRACING)
@ -132,3 +135,12 @@ KERNEL_STRUCT_MEMBER(guiding, float, sample_volume_guiding_rand, KERNEL_FEATURE_
/* The probability to use surface guiding (i.e., diffuse sampling prob * guiding prob). */
KERNEL_STRUCT_MEMBER(guiding, float, volume_guiding_sampling_prob, KERNEL_FEATURE_PATH_GUIDING)
KERNEL_STRUCT_END(guiding)
/******************************* Shadow linking *******************************/
KERNEL_STRUCT_BEGIN(shadow_link)
KERNEL_STRUCT_MEMBER(shadow_link, float, dedicated_light_weight, KERNEL_FEATURE_SHADOW_LINKING)
/* Copy of primitive and object from the last main path intersection. */
KERNEL_STRUCT_MEMBER(shadow_link, int, last_isect_prim, KERNEL_FEATURE_SHADOW_LINKING)
KERNEL_STRUCT_MEMBER(shadow_link, int, last_isect_object, KERNEL_FEATURE_SHADOW_LINKING)
KERNEL_STRUCT_END(shadow_link)

30
intern/cycles/kernel/integrator/state_util.h
View File

@ -60,6 +60,36 @@ ccl_device_forceinline void integrator_state_read_shadow_ray(ConstIntegratorShad
ray->dD = differential_zero_compact();
}
ccl_device_forceinline void integrator_state_write_shadow_ray_self(
KernelGlobals kg, IntegratorShadowState state, ccl_private const Ray *ccl_restrict ray)
{
if (kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING) {
INTEGRATOR_STATE_WRITE(state, shadow_ray, self_light) = ray->self.light;
}
/* Save memory by storing the light and object indices in the shadow_isect. */
/* TODO(sergey): This optimization does not work on GPU where multiple iterations of intersection
* is needed if there are more than 4 transparent intersections. The indices starts to conflict
* with each other. */
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, 0, object) = ray->self.object;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, 0, prim) = ray->self.prim;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, 1, object) = ray->self.light_object;
INTEGRATOR_STATE_ARRAY_WRITE(state, shadow_isect, 1, prim) = ray->self.light_prim;
}
ccl_device_forceinline void integrator_state_read_shadow_ray_self(
KernelGlobals kg, ConstIntegratorShadowState state, ccl_private Ray *ccl_restrict ray)
{
if (kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING) {
ray->self.light = INTEGRATOR_STATE(state, shadow_ray, self_light);
}
ray->self.object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, object);
ray->self.prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 0, prim);
ray->self.light_object = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, object);
ray->self.light_prim = INTEGRATOR_STATE_ARRAY(state, shadow_isect, 1, prim);
}
/* Intersection */
ccl_device_forceinline void integrator_state_write_isect(

3
intern/cycles/kernel/integrator/subsurface_disk.h
View File

@ -91,7 +91,8 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
/* Intersect with the same object. if multiple intersections are found it
* will use at most BSSRDF_MAX_HITS hits, a random subset of all hits. */

1
intern/cycles/kernel/integrator/subsurface_random_walk.h
View File

@ -197,6 +197,7 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
ray.self.prim = prim;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
/* Convert subsurface to volume coefficients.
* The single-scattering albedo is named alpha to avoid confusion with the surface albedo. */

20
intern/cycles/kernel/light/common.h
View File

@ -58,4 +58,24 @@ ccl_device float lamp_light_pdf(const float3 Ng, const float3 I, float t)
return t * t / cos_pi;
}
/* Visibility flag om the light shader. */
ccl_device_inline bool is_light_shader_visible_to_path(const int shader, const uint32_t path_flag)
{
if ((shader & SHADER_EXCLUDE_ANY) == 0) {
return true;
}
if (((shader & SHADER_EXCLUDE_DIFFUSE) && (path_flag & PATH_RAY_DIFFUSE)) ||
((shader & SHADER_EXCLUDE_GLOSSY) && ((path_flag & (PATH_RAY_GLOSSY | PATH_RAY_REFLECT)) ==
(PATH_RAY_GLOSSY | PATH_RAY_REFLECT))) ||
((shader & SHADER_EXCLUDE_TRANSMIT) && (path_flag & PATH_RAY_TRANSMIT)) ||
((shader & SHADER_EXCLUDE_CAMERA) && (path_flag & PATH_RAY_CAMERA)) ||
((shader & SHADER_EXCLUDE_SCATTER) && (path_flag & PATH_RAY_VOLUME_SCATTER)))
{
return false;
}
return true;
}
CCL_NAMESPACE_END

33
intern/cycles/kernel/light/distant.h
View File

@ -36,6 +36,39 @@ ccl_device_inline bool distant_light_sample(const ccl_global KernelLight *klight
return true;
}
/* Special intersection check.
* Returns true if the distant_light_sample_from_intersection() for this light would return true.
*
* The intersection parameters t, u, v are optimized for the shadow ray towards a dedicated light:
* u = v = 0, t = FLT_MAX.
*/
ccl_device bool distant_light_intersect(const ccl_global KernelLight *klight,
const ccl_private Ray *ccl_restrict ray,
ccl_private float *t,
ccl_private float *u,
ccl_private float *v)
{
kernel_assert(klight->type == LIGHT_DISTANT);
if (klight->distant.radius == 0.0f) {
return false;
}
const float3 lightD = klight->co;
const float costheta = dot(-lightD, ray->D);
const float cosangle = klight->distant.cosangle;
if (costheta < cosangle) {
return false;
}
*t = FLT_MAX;
*u = 0.0f;
*v = 0.0f;
return true;
}
ccl_device bool distant_light_sample_from_intersection(KernelGlobals kg,
const float3 ray_D,
const int lamp,

3
intern/cycles/kernel/light/distribution.h
View File

@ -48,6 +48,7 @@ ccl_device_noinline bool light_distribution_sample(KernelGlobals kg,
const float randv,
const float time,
const float3 P,
const int object_receiver,
const int bounce,
const uint32_t path_flag,
ccl_private LightSample *ls)
@ -56,7 +57,7 @@ ccl_device_noinline bool light_distribution_sample(KernelGlobals kg,
const int index = light_distribution_sample(kg, randn);
const float pdf_selection = kernel_data.integrator.distribution_pdf_lights;
return light_sample<in_volume_segment>(
kg, randu, randv, time, P, bounce, path_flag, index, 0, pdf_selection, ls);
kg, randu, randv, time, P, object_receiver, bounce, path_flag, index, 0, pdf_selection, ls);
}
ccl_device_inline float light_distribution_pdf_lamp(KernelGlobals kg)

242
intern/cycles/kernel/light/light.h
View File

@ -9,6 +9,7 @@
#include "kernel/light/point.h"
#include "kernel/light/spot.h"
#include "kernel/light/triangle.h"
#include "kernel/sample/lcg.h"
#include "kernel/sample/mapping.h"
@ -21,6 +22,72 @@ ccl_device_inline bool light_select_reached_max_bounces(KernelGlobals kg, int in
return (bounce > kernel_data_fetch(lights, index).max_bounces);
}
/* Light linking. */
ccl_device_inline int light_link_receiver_nee(KernelGlobals kg, const ccl_private ShaderData *sd)
{
#ifdef __LIGHT_LINKING__
if (!(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING)) {
return OBJECT_NONE;
}
return sd->object;
#else
return OBJECT_NONE;
#endif
}
ccl_device_inline int light_link_receiver_forward(KernelGlobals kg, IntegratorState state)
{
#ifdef __LIGHT_LINKING__
if (!(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING)) {
return OBJECT_NONE;
}
return INTEGRATOR_STATE(state, path, mis_ray_object);
#else
return OBJECT_NONE;
#endif
}
ccl_device_inline bool light_link_light_match(KernelGlobals kg,
const int object_receiver,
const int light_emitter)
{
#ifdef __LIGHT_LINKING__
if (!(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING)) {
return true;
}
const uint64_t set_membership = kernel_data_fetch(lights, light_emitter).light_set_membership;
const uint receiver_set = (object_receiver != OBJECT_NONE) ?
kernel_data_fetch(objects, object_receiver).receiver_light_set :
0;
return ((uint64_t(1) << uint64_t(receiver_set)) & set_membership) != 0;
#else
return true;
#endif
}
ccl_device_inline bool light_link_object_match(KernelGlobals kg,
const int object_receiver,
const int object_emitter)
{
#ifdef __LIGHT_LINKING__
if (!(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING)) {
return true;
}
const uint64_t set_membership = kernel_data_fetch(objects, object_emitter).light_set_membership;
const uint receiver_set = (object_receiver != OBJECT_NONE) ?
kernel_data_fetch(objects, object_receiver).receiver_light_set :
0;
return ((uint64_t(1) << uint64_t(receiver_set)) & set_membership) != 0;
#else
return true;
#endif
}
/* Sample point on an individual light. */
template<bool in_volume_segment>
@ -105,6 +172,7 @@ ccl_device_noinline bool light_sample(KernelGlobals kg,
const float randv,
const float time,
const float3 P,
const int object_receiver,
const int bounce,
const uint32_t path_flag,
const int emitter_index,
@ -138,6 +206,10 @@ ccl_device_noinline bool light_sample(KernelGlobals kg,
/* Mesh light. */
const int object = mesh_light.object_id;
if (!light_link_object_match(kg, object_receiver, object)) {
return false;
}
/* Exclude synthetic meshes from shadow catcher pass. */
if ((path_flag & PATH_RAY_SHADOW_CATCHER_PASS) &&
!(kernel_data_fetch(object_flag, object) & SD_OBJECT_SHADOW_CATCHER))
@ -152,11 +224,17 @@ ccl_device_noinline bool light_sample(KernelGlobals kg,
ls->shader |= shader_flag;
}
else {
if (UNLIKELY(light_select_reached_max_bounces(kg, ~prim, bounce))) {
const int light = ~prim;
if (!light_link_light_match(kg, object_receiver, light)) {
return false;
}
if (!light_sample<in_volume_segment>(kg, ~prim, randu, randv, P, path_flag, ls)) {
if (UNLIKELY(light_select_reached_max_bounces(kg, light, bounce))) {
return false;
}
if (!light_sample<in_volume_segment>(kg, light, randu, randv, P, path_flag, ls)) {
return false;
}
}
@ -167,15 +245,25 @@ ccl_device_noinline bool light_sample(KernelGlobals kg,
/* Intersect ray with individual light. */
ccl_device bool lights_intersect(KernelGlobals kg,
IntegratorState state,
ccl_private const Ray *ccl_restrict ray,
ccl_private Intersection *ccl_restrict isect,
const int last_prim,
const int last_object,
const int last_type,
const uint32_t path_flag)
/* Returns the total number of hits (the input num_hits plus the number of the new intersections).
*/
template<bool is_main_path>
ccl_device_forceinline int lights_intersect_impl(KernelGlobals kg,
ccl_private const Ray *ccl_restrict ray,
ccl_private Intersection *ccl_restrict isect,
const int last_prim,
const int last_object,
const int last_type,
const uint32_t path_flag,
const uint8_t path_mnee,
const int receiver_forward,
ccl_private uint *lcg_state,
int num_hits)
{
#ifdef __SHADOW_LINKING__
const bool is_indirect_ray = !(path_flag & PATH_RAY_CAMERA);
#endif
for (int lamp = 0; lamp < kernel_data.integrator.num_lights; lamp++) {
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
@ -192,9 +280,7 @@ ccl_device bool lights_intersect(KernelGlobals kg,
#ifdef __MNEE__
/* This path should have been resolved with mnee, it will
* generate a firefly for small lights since it is improbable. */
if ((INTEGRATOR_STATE(state, path, mnee) & PATH_MNEE_CULL_LIGHT_CONNECTION) &&
klight->use_caustics)
{
if ((path_mnee & PATH_MNEE_CULL_LIGHT_CONNECTION) && klight->use_caustics) {
continue;
}
#endif
@ -206,6 +292,32 @@ ccl_device bool lights_intersect(KernelGlobals kg,
}
}
#ifdef __SHADOW_LINKING__
/* For the main path exclude shadow-linked lights if intersecting with an indirect light ray.
* Those lights are handled via dedicated light intersect and shade kernels.
* For the shadow path used for the dedicated light shading ignore all non-shadow-linked
* lights. */
if (kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING) {
if (is_main_path) {
if (is_indirect_ray &&
kernel_data_fetch(lights, lamp).shadow_set_membership != LIGHT_LINK_MASK_ALL)
{
continue;
}
}
else if (kernel_data_fetch(lights, lamp).shadow_set_membership == LIGHT_LINK_MASK_ALL) {
continue;
}
}
#endif
#ifdef __LIGHT_LINKING__
/* Light linking. */
if (!light_link_light_match(kg, receiver_forward, lamp)) {
continue;
}
#endif
LightType type = (LightType)klight->type;
float t = 0.0f, u = 0.0f, v = 0.0f;
@ -224,25 +336,111 @@ ccl_device bool lights_intersect(KernelGlobals kg,
continue;
}
}
else if (type == LIGHT_DISTANT) {
if (is_main_path || ray->tmax != FLT_MAX) {
continue;
}
if (!distant_light_intersect(klight, ray, &t, &u, &v)) {
continue;
}
}
else {
continue;
}
if (t < isect->t &&
!(last_prim == lamp && last_object == OBJECT_NONE && last_type == PRIMITIVE_LAMP))
{
isect->t = t;
isect->u = u;
isect->v = v;
isect->type = PRIMITIVE_LAMP;
isect->prim = lamp;
isect->object = OBJECT_NONE;
/* Avoid self-intersections. */
if (last_prim == lamp && last_object == OBJECT_NONE && last_type == PRIMITIVE_LAMP) {
continue;
}
++num_hits;
#ifdef __SHADOW_LINKING__
if (!is_main_path) {
/* The non-main rays are only raced by the dedicated light kernel, after the shadow linking
* feature check. */
kernel_assert(kernel_data.kernel_features & KERNEL_FEATURE_SHADOW_LINKING);
if ((isect->prim != PRIM_NONE) && (lcg_step_float(lcg_state) > 1.0f / num_hits)) {
continue;
}
}
else
#endif
if (t >= isect->t)
{
continue;
}
isect->t = t;
isect->u = u;
isect->v = v;
isect->type = PRIMITIVE_LAMP;
isect->prim = lamp;
isect->object = OBJECT_NONE;
}
return num_hits;
}
/* Lights intersection for the main path.
* Intersects spot, point, and area lights. */
ccl_device bool lights_intersect(KernelGlobals kg,
IntegratorState state,
ccl_private const Ray *ccl_restrict ray,
ccl_private Intersection *ccl_restrict isect,
const int last_prim,
const int last_object,
const int last_type,
const uint32_t path_flag)
{
const uint8_t path_mnee = INTEGRATOR_STATE(state, path, mnee);
const int receiver_forward = light_link_receiver_forward(kg, state);
lights_intersect_impl<true>(kg,
ray,
isect,
last_prim,
last_object,
last_type,
path_flag,
path_mnee,
receiver_forward,
nullptr,
0);
return isect->prim != PRIM_NONE;
}
/* Lights intersection for the shadow linking.
* Intersects spot, point, area, and distant lights.
*
* Returns the total number of hits (the input num_hits plus the number of the new intersections).
*/
ccl_device int lights_intersect_shadow_linked(KernelGlobals kg,
ccl_private const Ray *ccl_restrict ray,
ccl_private Intersection *ccl_restrict isect,
const int last_prim,
const int last_object,
const int last_type,
const uint32_t path_flag,
const int receiver_forward,
ccl_private uint *lcg_state,
const int num_hits)
{
return lights_intersect_impl<false>(kg,
ray,
isect,
last_prim,
last_object,
last_type,
path_flag,
PATH_MNEE_NONE,
receiver_forward,
lcg_state,
num_hits);
}
/* Setup light sample from intersection. */
ccl_device bool light_sample_from_intersection(KernelGlobals kg,

84
intern/cycles/kernel/light/sample.h
View File

@ -250,6 +250,7 @@ ccl_device_inline void shadow_ray_setup(ccl_private const ShaderData *ccl_restri
ray->self.prim = (skip_self) ? sd->prim : PRIM_NONE;
ray->self.light_object = ls->object;
ray->self.light_prim = ls->prim;
ray->self.light = ls->lamp;
}
/* Create shadow ray towards light sample. */
@ -324,20 +325,32 @@ ccl_device_inline bool light_sample_from_volume_segment(KernelGlobals kg,
const float3 P,
const float3 D,
const float t,
const int object_receiver,
const int bounce,
const uint32_t path_flag,
ccl_private LightSample *ls)
{
#ifdef __LIGHT_TREE__
if (kernel_data.integrator.use_light_tree) {
return light_tree_sample<true>(
kg, randn, randu, randv, time, P, D, t, SD_BSDF_HAS_TRANSMISSION, bounce, path_flag, ls);
return light_tree_sample<true>(kg,
randn,
randu,
randv,
time,
P,
D,
t,
object_receiver,
SD_BSDF_HAS_TRANSMISSION,
bounce,
path_flag,
ls);
}
else
#endif
{
return light_distribution_sample<true>(
kg, randn, randu, randv, time, P, bounce, path_flag, ls);
kg, randn, randu, randv, time, P, object_receiver, bounce, path_flag, ls);
}
}
@ -349,6 +362,7 @@ ccl_device bool light_sample_from_position(KernelGlobals kg,
const float time,
const float3 P,
const float3 N,
const int object_receiver,
const int shader_flags,
const int bounce,
const uint32_t path_flag,
@ -356,47 +370,25 @@ ccl_device bool light_sample_from_position(KernelGlobals kg,
{
#ifdef __LIGHT_TREE__
if (kernel_data.integrator.use_light_tree) {
return light_tree_sample<false>(
kg, randn, randu, randv, time, P, N, 0, shader_flags, bounce, path_flag, ls);
return light_tree_sample<false>(kg,
randn,
randu,
randv,
time,
P,
N,
0.0f,
object_receiver,
shader_flags,
bounce,
path_flag,
ls);
}
else
#endif
{
return light_distribution_sample<false>(
kg, randn, randu, randv, time, P, bounce, path_flag, ls);
}
}
ccl_device_inline bool light_sample_new_position(KernelGlobals kg,
const float randu,
const float randv,
const float time,
const float3 P,
ccl_private LightSample *ls)
{
/* Sample a new position on the same light, for volume sampling. */
if (ls->type == LIGHT_TRIANGLE) {
if (!triangle_light_sample<false>(kg, ls->prim, ls->object, randu, randv, time, ls, P)) {
return false;
}
#ifdef __LIGHT_TREE__
if (kernel_data.integrator.use_light_tree) {
ls->pdf *= ls->pdf_selection;
}
else
#endif
{
/* Handled in triangle_light_sample for efficiency. */
}
return true;
}
else {
if (!light_sample<false>(kg, ls->lamp, randu, randv, P, 0, ls)) {
return false;
}
ls->pdf *= ls->pdf_selection;
return true;
kg, randn, randu, randv, time, P, object_receiver, bounce, path_flag, ls);
}
}
@ -440,7 +432,8 @@ ccl_device_inline float light_sample_mis_weight_forward_surface(KernelGlobals kg
uint prim_offset = kernel_data_fetch(object_prim_offset, sd->object);
uint triangle = kernel_data_fetch(triangle_to_tree, sd->prim - prim_offset + lookup_offset);
pdf *= light_tree_pdf(kg, ray_P, N, path_flag, sd->object, triangle);
pdf *= light_tree_pdf(
kg, ray_P, N, path_flag, sd->object, triangle, light_link_receiver_forward(kg, state));
}
else
#endif
@ -464,7 +457,13 @@ ccl_device_inline float light_sample_mis_weight_forward_lamp(KernelGlobals kg,
#ifdef __LIGHT_TREE__
if (kernel_data.integrator.use_light_tree) {
const float3 N = INTEGRATOR_STATE(state, path, mis_origin_n);
pdf *= light_tree_pdf(kg, P, N, path_flag, 0, kernel_data_fetch(light_to_tree, ls->lamp));
pdf *= light_tree_pdf(kg,
P,
N,
path_flag,
0,
kernel_data_fetch(light_to_tree, ls->lamp),
light_link_receiver_forward(kg, state));
}
else
#endif
@ -499,7 +498,8 @@ ccl_device_inline float light_sample_mis_weight_forward_background(KernelGlobals
if (kernel_data.integrator.use_light_tree) {
const float3 N = INTEGRATOR_STATE(state, path, mis_origin_n);
uint light = kernel_data_fetch(light_to_tree, kernel_data.background.light_index);
pdf *= light_tree_pdf(kg, ray_P, N, path_flag, 0, light);
pdf *= light_tree_pdf(
kg, ray_P, N, path_flag, 0, light, light_link_receiver_forward(kg, state));
}
else
#endif

70
intern/cycles/kernel/light/tree.h
View File

@ -583,7 +583,7 @@ ccl_device int light_tree_cluster_select_emitter(KernelGlobals kg,
else {
selected_index = -1;
for (int i = 0; i < knode->num_emitters; i++) {
int current_index = knode->inner.right_child + i;
int current_index = knode->leaf.first_emitter + i;
sample_resevoir(current_index,
float(has_importance & 1),
selected_index,
@ -656,6 +656,19 @@ ccl_device bool get_left_probability(KernelGlobals kg,
return true;
}
ccl_device int light_tree_root_node_index(KernelGlobals kg, const int object_receiver)
{
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_LINKING) {
const uint receiver_light_set =
(object_receiver != OBJECT_NONE) ?
kernel_data_fetch(objects, object_receiver).receiver_light_set :
0;
return kernel_data.light_link_sets[receiver_light_set].light_tree_root;
}
return 0;
}
template<bool in_volume_segment>
ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
float randn,
@ -665,6 +678,7 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
const float3 P,
float3 N_or_D,
float t,
const int object_receiver,
const int shader_flags,
const int bounce,
const uint32_t path_flag,
@ -678,8 +692,8 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
float pdf_leaf = 1.0f;
float pdf_selection = 1.0f;
int selected_emitter = -1;
int object = 0;
int node_index = 0; /* Root node. */
int object_emitter = 0;
int node_index = light_tree_root_node_index(kg, object_receiver);
float3 local_P = P;
@ -697,14 +711,13 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
}
else {
/* Continue with the picked mesh light. */
object = kernel_data_fetch(light_tree_emitters, selected_emitter).mesh.object_id;
object_emitter = kernel_data_fetch(light_tree_emitters, selected_emitter).mesh.object_id;
continue;
}
}
/* At an interior node, the left child is directly after the parent, while the right child is
* stored as the child index. */
const int left_index = node_index + 1;
/* Inner node. */
const int left_index = knode->inner.left_child;
const int right_index = knode->inner.right_child;
float left_prob;
@ -727,25 +740,40 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
pdf_selection *= pdf_leaf;
return light_sample<in_volume_segment>(
kg, randu, randv, time, P, bounce, path_flag, selected_emitter, object, pdf_selection, ls);
return light_sample<in_volume_segment>(kg,
randu,
randv,
time,
P,
object_receiver,
bounce,
path_flag,
selected_emitter,
object_emitter,
pdf_selection,
ls);
}
/* We need to be able to find the probability of selecting a given light for MIS. */
ccl_device float light_tree_pdf(
KernelGlobals kg, float3 P, float3 N, const int path_flag, const int object, const uint target)
ccl_device float light_tree_pdf(KernelGlobals kg,
float3 P,
float3 N,
const int path_flag,
const int object_emitter,
const uint index_emitter,
const int object_receiver)
{
const bool has_transmission = (path_flag & PATH_RAY_MIS_HAD_TRANSMISSION);
ccl_global const KernelLightTreeEmitter *kemitter = &kernel_data_fetch(light_tree_emitters,
target);
index_emitter);
int root_index;
uint bit_trail, target_emitter;
if (is_triangle(kemitter)) {
/* If the target is an emissive triangle, first traverse the top level tree to find the mesh
* light emitter, then traverse the subtree. */
target_emitter = kernel_data_fetch(object_to_tree, object);
target_emitter = kernel_data_fetch(object_to_tree, object_emitter);
ccl_global const KernelLightTreeEmitter *kmesh = &kernel_data_fetch(light_tree_emitters,
target_emitter);
root_index = kmesh->mesh.node_id;
@ -759,11 +787,11 @@ ccl_device float light_tree_pdf(
else {
root_index = 0;
bit_trail = kemitter->bit_trail;
target_emitter = target;
target_emitter = index_emitter;
}
float pdf = 1.0f;
int node_index = 0;
int node_index = light_tree_root_node_index(kg, object_receiver);
/* Traverse the light tree until we reach the target leaf node. */
while (true) {
@ -802,11 +830,11 @@ ccl_device float light_tree_pdf(
if (root_index) {
/* Arrived at the mesh light. Continue with the subtree. */
float unused;
light_tree_to_local_space<false>(kg, object, P, N, unused);
light_tree_to_local_space<false>(kg, object_emitter, P, N, unused);
node_index = root_index;
root_index = 0;
target_emitter = target;
target_emitter = index_emitter;
bit_trail = kemitter->bit_trail;
continue;
}
@ -815,8 +843,8 @@ ccl_device float light_tree_pdf(
}
}
/* Interior node. */
const int left_index = node_index + 1;
/* Inner node. */
const int left_index = knode->inner.left_child;
const int right_index = knode->inner.right_child;
float left_prob;
@ -826,10 +854,12 @@ ccl_device float light_tree_pdf(
return 0.0f;
}
bit_trail >>= kernel_data_fetch(light_tree_nodes, node_index).bit_skip;
const bool go_left = (bit_trail & 1) == 0;
bit_trail >>= 1;
pdf *= go_left ? left_prob : (1.0f - left_prob);
node_index = go_left ? left_index : right_index;
pdf *= go_left ? left_prob : (1.0f - left_prob);
if (pdf == 0) {
return 0.0f;

1
intern/cycles/kernel/osl/services.cpp
View File

@ -1682,6 +1682,7 @@ bool OSLRenderServices::trace(TraceOpt &options,
ray.self.prim = PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
if (options.mindist == 0.0f) {
/* avoid self-intersections */

1
intern/cycles/kernel/svm/ao.h
View File

@ -66,6 +66,7 @@ ccl_device float svm_ao(
ray.self.prim = sd->prim;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
ray.dP = differential_zero_compact();
ray.dD = differential_zero_compact();

1
intern/cycles/kernel/svm/bevel.h
View File

@ -188,6 +188,7 @@ ccl_device float3 svm_bevel(
ray.self.prim = PRIM_NONE;
ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE;
ray.self.light = LAMP_NONE;
/* Intersect with the same object. if multiple intersections are found it
* will use at most LOCAL_MAX_HITS hits, a random subset of all hits. */

37
intern/cycles/kernel/types.h
View File

@ -48,6 +48,9 @@ CCL_NAMESPACE_BEGIN
#define PASS_UNUSED (~0)
#define LIGHTGROUP_NONE (~0)
#define LIGHT_LINK_SET_MAX 64
#define LIGHT_LINK_MASK_ALL (~uint64_t(0))
#define INTEGRATOR_SHADOW_ISECT_SIZE_CPU 1024U
#define INTEGRATOR_SHADOW_ISECT_SIZE_GPU 4U
@ -64,6 +67,8 @@ CCL_NAMESPACE_BEGIN
#define __DENOISING_FEATURES__
#define __DPDU__
#define __HAIR__
#define __LIGHT_LINKING__
#define __SHADOW_LINKING__
#define __LIGHT_TREE__
#define __OBJECT_MOTION__
#define __PASSES__
@ -313,6 +318,8 @@ enum PathRayFlag : uint32_t {
// 8bit enum, just in case we need to move more variables in it
enum PathRayMNEE {
PATH_MNEE_NONE = 0,
PATH_MNEE_VALID = (1U << 0U),
PATH_MNEE_RECEIVER_ANCESTOR = (1U << 1U),
PATH_MNEE_CULL_LIGHT_CONNECTION = (1U << 2U),
@ -579,6 +586,7 @@ typedef struct RaySelfPrimitives {
int object; /* Instance prim is a part of */
int light_prim; /* Light primitive */
int light_object; /* Light object */
int light; /* Light ID (the light the shadow ray is traced towards to) */
} RaySelfPrimitives;
typedef struct Ray {
@ -1228,6 +1236,10 @@ typedef struct KernelBake {
} KernelBake;
static_assert_align(KernelBake, 16);
typedef struct KernelLightLinkSet {
uint light_tree_root;
} KernelLightLinkSet;
typedef struct KernelData {
/* Features and limits. */
uint kernel_features;
@ -1239,6 +1251,7 @@ typedef struct KernelData {
KernelCamera cam;
KernelBake bake;
KernelTables tables;
KernelLightLinkSet light_link_sets[LIGHT_LINK_SET_MAX];
/* Potentially specialized data members. */
#define KERNEL_STRUCT_BEGIN(name, parent) name parent;
@ -1304,6 +1317,12 @@ typedef struct KernelObject {
/* Volume velocity scale. */
float velocity_scale;
/* TODO: separate array to avoid memory overhead when not used. */
uint64_t light_set_membership;
uint receiver_light_set;
uint64_t shadow_set_membership;
uint blocker_shadow_set;
} KernelObject;
static_assert_align(KernelObject, 16);
@ -1369,6 +1388,8 @@ typedef struct KernelLight {
KernelAreaLight area;
KernelDistantLight distant;
};
uint64_t light_set_membership;
uint64_t shadow_set_membership;
} KernelLight;
static_assert_align(KernelLight, 16);
@ -1423,7 +1444,9 @@ typedef struct KernelLightTreeNode {
int first_emitter; /* The index of the first emitter. */
} leaf;
struct {
int right_child; /* The index of the right child. */
/* Indices of the children. */
int left_child;
int right_child;
} inner;
struct {
int reference; /* A reference to the node with the subtree. */
@ -1432,6 +1455,12 @@ typedef struct KernelLightTreeNode {
/* Bit trail. */
uint bit_trail;
/* Bits to skip in the bit trail, to skip nodes in for specialized trees. */
uint8_t bit_skip;
/* Padding. */
uint8_t pad[11];
} KernelLightTreeNode;
static_assert_align(KernelLightTreeNode, 16);
@ -1554,6 +1583,7 @@ typedef enum DeviceKernel : int {
DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT,
DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND,
DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE,
@ -1561,6 +1591,7 @@ typedef enum DeviceKernel : int {
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE,
DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW,
DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT,
DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL,
DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY,
@ -1680,6 +1711,10 @@ enum KernelFeatureFlag : uint32_t {
/* OSL. */
KERNEL_FEATURE_OSL = (1U << 26U),
/* Light and shadow linking. */
KERNEL_FEATURE_LIGHT_LINKING = (1U << 27U),
KERNEL_FEATURE_SHADOW_LINKING = (1U << 28U),
};
/* Shader node feature mask, to specialize shader evaluation for kernels. */

216
intern/cycles/scene/light.cpp
View File

@ -138,6 +138,8 @@ NODE_DEFINE(Light)
SOCKET_NODE(shader, "Shader", Shader::get_node_type());
SOCKET_STRING(lightgroup, "Light Group", ustring());
SOCKET_UINT64(light_set_membership, "Light Set Membership", LIGHT_LINK_MASK_ALL);
SOCKET_UINT64(shadow_set_membership, "Shadow Set Membership", LIGHT_LINK_MASK_ALL);
SOCKET_BOOLEAN(normalize, "Normalize", true);
@ -448,7 +450,8 @@ struct LightTreeFlatten {
static void light_tree_node_copy_to_device(KernelLightTreeNode &knode,
const LightTreeNode &node,
const int child_index)
const int left_child,
const int right_child)
{
/* Convert node to kernel representation. */
knode.energy = node.measure.energy;
@ -461,6 +464,7 @@ static void light_tree_node_copy_to_device(KernelLightTreeNode &knode,
knode.bcone.theta_e = node.measure.bcone.theta_e;
knode.bit_trail = node.bit_trail;
knode.bit_skip = 0;
knode.type = static_cast<LightTreeNodeType>(node.type);
if (node.is_leaf() || node.is_distant()) {
@ -469,7 +473,8 @@ static void light_tree_node_copy_to_device(KernelLightTreeNode &knode,
}
else if (node.is_inner()) {
knode.num_emitters = -1;
knode.inner.right_child = child_index;
knode.inner.left_child = left_child;
knode.inner.right_child = right_child;
}
}
@ -584,21 +589,18 @@ static int light_tree_flatten(LightTreeFlatten &flatten,
{
/* Convert both inner nodes and primitives to device representation. */
const int node_index = next_node_index++;
int child_index = -1;
int left_child = -1, right_child = -1;
if (node->is_leaf() || node->is_distant()) {
light_tree_leaf_emitters_copy_and_flatten(flatten, *node, knodes, kemitters, next_node_index);
}
else if (node->is_inner()) {
/* Nodes are stored in depth first order so that the left child node
* immediately follows the parent, and only the right child index needs
* to be stored. */
light_tree_flatten(flatten,
node->get_inner().children[LightTree::left].get(),
knodes,
kemitters,
next_node_index);
child_index = light_tree_flatten(flatten,
left_child = light_tree_flatten(flatten,
node->get_inner().children[LightTree::left].get(),
knodes,
kemitters,
next_node_index);
right_child = light_tree_flatten(flatten,
node->get_inner().children[LightTree::right].get(),
knodes,
kemitters,
@ -609,11 +611,137 @@ static int light_tree_flatten(LightTreeFlatten &flatten,
assert(node->is_instance());
}
light_tree_node_copy_to_device(knodes[node_index], *node, child_index);
light_tree_node_copy_to_device(knodes[node_index], *node, left_child, right_child);
return node_index;
}
static void light_tree_emitters_copy_and_flatten(LightTreeFlatten &flatten,
const LightTreeNode *node,
KernelLightTreeNode *knodes,
KernelLightTreeEmitter *kemitters,
int &next_node_index)
{
/* Convert only emitters to device representation. */
if (node->is_leaf() || node->is_distant()) {
light_tree_leaf_emitters_copy_and_flatten(flatten, *node, knodes, kemitters, next_node_index);
}
else {
assert(node->is_inner());
light_tree_emitters_copy_and_flatten(flatten,
node->get_inner().children[LightTree::left].get(),
knodes,
kemitters,
next_node_index);
light_tree_emitters_copy_and_flatten(flatten,
node->get_inner().children[LightTree::right].get(),
knodes,
kemitters,
next_node_index);
}
}
static std::pair<int, LightTreeMeasure> light_tree_specialize_nodes_flatten(
const LightTreeFlatten &flatten,
LightTreeNode *node,
const uint64_t light_link_mask,
const int depth,
vector<KernelLightTreeNode> &knodes,
int &next_node_index)
{
assert(!node->is_instance());
/* Convert inner nodes to device representation, specialized for light linking. */
int node_index, left_child = -1, right_child = -1;
LightTreeNode new_node(LightTreeMeasure::empty, node->bit_trail);
if (depth == 0 && !(node->light_link.set_membership & light_link_mask)) {
/* Ensure there is always a root node. */
node_index = next_node_index++;
new_node.make_leaf(-1, 0);
}
else if (node->light_link.shareable && node->light_link.shared_node_index != -1) {
/* Share subtree already built for another light link set. */
return std::make_pair(node->light_link.shared_node_index, node->measure);
}
else if (node->is_leaf() || node->is_distant()) {
/* Specialize leaf node. */
node_index = next_node_index++;
int first_emitter = -1;
int num_emitters = 0;
for (int i = 0; i < node->get_leaf().num_emitters; i++) {
const LightTreeEmitter &emitter = flatten.emitters[node->get_leaf().first_emitter_index + i];
if (emitter.light_set_membership & light_link_mask) {
/* Assumes emitters are consecutive due to LighTree::sort_leaf. */
if (first_emitter == -1) {
first_emitter = node->get_leaf().first_emitter_index + i;
}
num_emitters++;
new_node.measure.add(emitter.measure);
}
}
assert(first_emitter != -1);
new_node.make_leaf(first_emitter, num_emitters);
}
else {
assert(node->is_inner());
assert(new_node.is_inner());
/* Specialize inner node. */
LightTreeNode *left_node = node->get_inner().children[LightTree::left].get();
LightTreeNode *right_node = node->get_inner().children[LightTree::right].get();
/* Skip nodes that have only one child. We have a single bit trail for each
* primitive, bit_skip is incremented in the child node to skip the bit for
* this parent node. */
LightTreeNode *only_node = nullptr;
if (!(left_node->light_link.set_membership & light_link_mask)) {
only_node = right_node;
}
else if (!(right_node->light_link.set_membership & light_link_mask)) {
only_node = left_node;
}
if (only_node) {
const auto [only_index, only_measure] = light_tree_specialize_nodes_flatten(
flatten, only_node, light_link_mask, depth + 1, knodes, next_node_index);
assert(only_index != -1);
knodes[only_index].bit_skip++;
return std::make_pair(only_index, only_measure);
}
/* Create inner node. */
node_index = next_node_index++;
const auto [left_index, left_measure] = light_tree_specialize_nodes_flatten(
flatten, left_node, light_link_mask, depth + 1, knodes, next_node_index);
const auto [right_index, right_measure] = light_tree_specialize_nodes_flatten(
flatten, right_node, light_link_mask, depth + 1, knodes, next_node_index);
new_node.measure = left_measure;
new_node.measure.add(right_measure);
left_child = left_index;
right_child = right_index;
}
/* Convert to kernel node. */
if (knodes.size() <= node_index) {
knodes.resize(node_index + 1);
}
light_tree_node_copy_to_device(knodes[node_index], new_node, left_child, right_child);
if (node->light_link.shareable) {
node->light_link.shared_node_index = node_index;
}
return std::make_pair(node_index, new_node.measure);
}
void LightManager::device_update_tree(Device *,
DeviceScene *dscene,
Scene *scene,
@ -647,23 +775,68 @@ void LightManager::device_update_tree(Device *,
flatten.mesh_array = dscene->object_to_tree.alloc(scene->objects.size());
flatten.triangle_array = dscene->triangle_to_tree.alloc(light_tree.num_triangles);
/* Allocate emitters and nodes. */
/* Allocate emitters */
const size_t num_emitters = light_tree.num_emitters();
KernelLightTreeEmitter *kemitters = dscene->light_tree_emitters.alloc(num_emitters);
KernelLightTreeNode *knodes = dscene->light_tree_nodes.alloc(light_tree.num_nodes);
/* Update integrator state. */
kintegrator->use_direct_light = num_emitters > 0;
/* Copy nodes and emitters. */
if (root) {
int next_node_index = 0;
light_tree_flatten(flatten, root, knodes, kemitters, next_node_index);
}
/* Test if light linking is used. */
const bool use_light_linking = root && (light_tree.light_link_receiver_used != 1);
KernelLightLinkSet *klight_link_sets = dscene->data.light_link_sets;
memset(klight_link_sets, 0, sizeof(dscene->data.light_link_sets));
VLOG_INFO << "Use light tree with " << num_emitters << " emitters and " << light_tree.num_nodes
<< " nodes.";
if (!use_light_linking) {
/* Regular light tree without linking. */
KernelLightTreeNode *knodes = dscene->light_tree_nodes.alloc(light_tree.num_nodes);
if (root) {
int next_node_index = 0;
light_tree_flatten(flatten, root, knodes, kemitters, next_node_index);
}
}
else {
int next_node_index = 0;
vector<KernelLightTreeNode> light_link_nodes;
/* Write primitives, and any subtrees for instances. */
if (root) {
/* Reserve enough size of all instance subtrees, then shrink back to
* actual number of nodes used. */
light_link_nodes.resize(light_tree.num_nodes);
light_tree_emitters_copy_and_flatten(
flatten, root, light_link_nodes.data(), kemitters, next_node_index);
light_link_nodes.resize(next_node_index);
}
/* Specialized light trees for linking. */
for (uint64_t tree_index = 0; tree_index < LIGHT_LINK_SET_MAX; tree_index++) {
const uint64_t tree_mask = uint64_t(1) << tree_index;
if (!(light_tree.light_link_receiver_used & tree_mask)) {
continue;
}
if (root) {
klight_link_sets[tree_index].light_tree_root =
light_tree_specialize_nodes_flatten(
flatten, root, tree_mask, 0, light_link_nodes, next_node_index)
.first;
}
}
/* Allocate and copy nodes into device array. */
KernelLightTreeNode *knodes = dscene->light_tree_nodes.alloc(light_link_nodes.size());
memcpy(knodes, light_link_nodes.data(), light_link_nodes.size() * sizeof(*knodes));
VLOG_INFO << "Specialized light tree for light linking, with "
<< light_link_nodes.size() - light_tree.num_nodes << " additional nodes.";
}
/* Copy arrays to device. */
dscene->light_tree_nodes.copy_to_device();
dscene->light_tree_emitters.copy_to_device();
@ -1152,6 +1325,9 @@ void LightManager::device_update_lights(Device *device, DeviceScene *dscene, Sce
klights[light_index].lightgroup = LIGHTGROUP_NONE;
}
klights[light_index].light_set_membership = light->light_set_membership;
klights[light_index].shadow_set_membership = light->shadow_set_membership;
light_index++;
}

2
intern/cycles/scene/light.h
View File

@ -72,6 +72,8 @@ class Light : public Node {
NODE_SOCKET_API(uint, random_id)
NODE_SOCKET_API(ustring, lightgroup)
NODE_SOCKET_API(uint64_t, light_set_membership);
NODE_SOCKET_API(uint64_t, shadow_set_membership);
NODE_SOCKET_API(bool, normalize)

44
intern/cycles/scene/light_tree.cpp
View File

@ -77,6 +77,7 @@ OrientationBounds merge(const OrientationBounds &cone_a, const OrientationBounds
LightTreeEmitter::LightTreeEmitter(Object *object, int object_id) : object_id(object_id)
{
centroid = object->bounds.center();
light_set_membership = object->get_light_set_membership();
}
LightTreeEmitter::LightTreeEmitter(Scene *scene,
@ -138,6 +139,8 @@ LightTreeEmitter::LightTreeEmitter(Scene *scene,
for (int i = 0; i < 3; i++) {
measure.bbox.grow(vertices[i]);
}
light_set_membership = object->get_light_set_membership();
}
else {
assert(is_light());
@ -216,6 +219,20 @@ LightTreeEmitter::LightTreeEmitter(Scene *scene,
/* Use absolute value of energy so lights with negative strength are properly supported in the
* light tree. */
measure.energy = fabsf(average(strength));
light_set_membership = lamp->get_light_set_membership();
}
}
static void sort_leaf(const int start, const int end, LightTreeEmitter *emitters)
{
/* Sort primitive by light link mask so that specialized trees can use a subset of these. */
if (end > start) {
std::sort(emitters + start,
emitters + end,
[](const LightTreeEmitter &a, const LightTreeEmitter &b) {
return a.light_set_membership < b.light_set_membership;
});
}
}
@ -279,6 +296,8 @@ LightTree::LightTree(Scene *scene,
return;
}
light_link_receiver_used |= (uint64_t(1) << object->get_receiver_light_set());
if (!object->usable_as_light()) {
object_id++;
continue;
@ -390,7 +409,15 @@ LightTreeNode *LightTree::build(Scene *scene, DeviceScene *dscene)
for (int i = 0; i < num_distant_lights; i++) {
root_->get_inner().children[right]->add(distant_lights_[i]);
}
sort_leaf(0, num_distant_lights, distant_lights_.data());
root_->get_inner().children[right]->make_distant(num_local_lights, num_distant_lights);
root_->measure = root_->get_inner().children[left]->measure +
root_->get_inner().children[right]->measure;
root_->light_link = root_->get_inner().children[left]->light_link +
root_->get_inner().children[right]->light_link;
std::move(distant_lights_.begin(), distant_lights_.end(), std::back_inserter(emitters_));
return root_.get();
@ -421,7 +448,7 @@ void LightTree::recursive_build(const Child child,
/* Find the best place to split the emitters into 2 nodes.
* If the best split cost is no better than making a leaf node, make a leaf instead. */
int split_dim = -1, middle;
if (should_split(emitters, start, middle, end, node->measure, split_dim)) {
if (should_split(emitters, start, middle, end, node->measure, node->light_link, split_dim)) {
if (split_dim != -1) {
/* Partition the emitters between start and end based on the centroids. */
@ -453,6 +480,7 @@ void LightTree::recursive_build(const Child child,
}
}
else {
sort_leaf(start, end, emitters);
node->make_leaf(start, end - start);
}
}
@ -462,13 +490,15 @@ bool LightTree::should_split(LightTreeEmitter *emitters,
int &middle,
const int end,
LightTreeMeasure &measure,
LightTreeLightLink &light_link,
int &split_dim)
{
const int num_emitters = end - start;
if (num_emitters < 2) {
if (num_emitters) {
/* Do not try to split if there is only one emitter. */
measure = (emitters + start)->measure;
measure = emitters[start].measure;
light_link = LightTreeLightLink(emitters[start].light_set_membership);
}
return false;
}
@ -519,6 +549,7 @@ bool LightTree::should_split(LightTreeEmitter *emitters,
if (dim == 0) {
/* Calculate node measure by summing up the bucket measure. */
measure = left_buckets.back().measure + buckets.back().measure;
light_link = left_buckets.back().light_link + buckets.back().light_link;
/* Degenerate case with co-located emitters. */
if (is_zero(centroid_bbox.size())) {
@ -569,7 +600,14 @@ __forceinline LightTreeMeasure operator+(const LightTreeMeasure &a, const LightT
LightTreeBucket operator+(const LightTreeBucket &a, const LightTreeBucket &b)
{
return LightTreeBucket(a.measure + b.measure, a.count + b.count);
return LightTreeBucket(a.measure + b.measure, a.light_link + b.light_link, a.count + b.count);
}
LightTreeLightLink operator+(const LightTreeLightLink &a, const LightTreeLightLink &b)
{
LightTreeLightLink c(a);
c.add(b);
return c;
}
CCL_NAMESPACE_END

62
intern/cycles/scene/light_tree.h
View File

@ -131,6 +131,48 @@ LightTreeMeasure operator+(const LightTreeMeasure &a, const LightTreeMeasure &b)
struct LightTreeNode;
/* Light Linking. */
struct LightTreeLightLink {
/* Bitmask for membership of primitives in this node. */
uint64_t set_membership = 0;
/* When all primitives below this node have identical light set membership, this
* part of the light tree can be shared between specialized trees. */
bool shareable = true;
int shared_node_index = -1;
LightTreeLightLink() = default;
LightTreeLightLink(const uint64_t set_membership) : set_membership(set_membership) {}
void add(const uint64_t prim_set_membership)
{
if (set_membership == 0) {
set_membership = prim_set_membership;
}
else if (prim_set_membership != set_membership) {
set_membership |= prim_set_membership;
shareable = false;
}
}
void add(const LightTreeLightLink &other)
{
if (set_membership == 0) {
set_membership = other.set_membership;
shareable = other.shareable;
}
else if (other.set_membership != set_membership) {
set_membership |= other.set_membership;
shareable = false;
}
else if (!other.shareable) {
shareable = false;
}
}
};
LightTreeLightLink operator+(const LightTreeLightLink &a, const LightTreeLightLink &b);
/* Light Tree Emitter
* An emitter is a built-in light, an emissive mesh, or an emissive triangle. */
struct LightTreeEmitter {
@ -144,6 +186,7 @@ struct LightTreeEmitter {
int object_id;
float3 centroid;
uint64_t light_set_membership;
LightTreeMeasure measure;
@ -170,19 +213,23 @@ struct LightTreeEmitter {
* Struct used to determine splitting costs in the light BVH. */
struct LightTreeBucket {
LightTreeMeasure measure;
LightTreeLightLink light_link;
int count = 0;
static const int num_buckets = 12;
LightTreeBucket() = default;
LightTreeBucket(const LightTreeMeasure &measure, const int &count)
: measure(measure), count(count)
LightTreeBucket(const LightTreeMeasure &measure,
const LightTreeLightLink &light_link,
const int &count)
: measure(measure), light_link(light_link), count(count)
{
}
void add(const LightTreeEmitter &emitter)
{
measure.add(emitter.measure);
light_link.add(emitter.light_set_membership);
count++;
}
};
@ -192,6 +239,7 @@ LightTreeBucket operator+(const LightTreeBucket &a, const LightTreeBucket &b);
/* Light Tree Node */
struct LightTreeNode {
LightTreeMeasure measure;
LightTreeLightLink light_link;
uint bit_trail;
int object_id;
@ -260,7 +308,7 @@ struct LightTreeNode {
return std::get<Instance>(variant_type);
}
void make_leaf(const int &first_emitter_index, const int &num_emitters)
void make_leaf(const int first_emitter_index, const int num_emitters)
{
variant_type = Leaf();
Leaf &leaf = get_leaf();
@ -270,7 +318,7 @@ struct LightTreeNode {
type = LIGHT_TREE_LEAF;
}
void make_distant(const int &first_emitter_index, const int &num_emitters)
void make_distant(const int first_emitter_index, const int num_emitters)
{
variant_type = Leaf();
Leaf &leaf = get_leaf();
@ -280,7 +328,7 @@ struct LightTreeNode {
type = LIGHT_TREE_DISTANT;
}
void make_instance(LightTreeNode *reference, const int &object_id)
void make_instance(LightTreeNode *reference, const int object_id)
{
variant_type = Instance();
Instance &instance = get_instance();
@ -344,6 +392,9 @@ class LightTree {
std::atomic<int> num_nodes = 0;
size_t num_triangles = 0;
/* Bitmask of receiver light sets used. Default set is always used. */
uint64_t light_link_receiver_used = 1;
/* An inner node itself or its left and right child. */
enum Child {
self = -1,
@ -392,6 +443,7 @@ class LightTree {
int &middle,
const int end,
LightTreeMeasure &measure,
LightTreeLightLink &light_link,
int &split_dim);
/* Check whether the light tree can use this triangle as light-emissive. */

12
intern/cycles/scene/object.cpp
View File

@ -100,6 +100,10 @@ NODE_DEFINE(Object)
SOCKET_FLOAT(ao_distance, "AO Distance", 0.0f);
SOCKET_STRING(lightgroup, "Light Group", ustring());
SOCKET_UINT(receiver_light_set, "Light Set Index", 0);
SOCKET_UINT64(light_set_membership, "Light Set Membership", LIGHT_LINK_MASK_ALL);
SOCKET_UINT(blocker_shadow_set, "Shadow Set Index", 0);
SOCKET_UINT64(shadow_set_membership, "Shadow Set Membership", LIGHT_LINK_MASK_ALL);
return type;
}
@ -456,6 +460,14 @@ void ObjectManager::device_update_object_transform(UpdateObjectTransformState *s
kobject.particle_index = particle_index;
kobject.motion_offset = 0;
kobject.ao_distance = ob->ao_distance;
kobject.receiver_light_set = ob->receiver_light_set >= LIGHT_LINK_SET_MAX ?
0 :
ob->receiver_light_set;
kobject.light_set_membership = ob->light_set_membership;
kobject.blocker_shadow_set = ob->blocker_shadow_set >= LIGHT_LINK_SET_MAX ?
0 :
ob->blocker_shadow_set;
kobject.shadow_set_membership = ob->shadow_set_membership;
if (geom->get_use_motion_blur()) {
state->have_motion = true;

4
intern/cycles/scene/object.h
View File

@ -68,6 +68,10 @@ class Object : public Node {
NODE_SOCKET_API(float, ao_distance)
NODE_SOCKET_API(ustring, lightgroup)
NODE_SOCKET_API(uint, receiver_light_set)
NODE_SOCKET_API(uint64_t, light_set_membership)
NODE_SOCKET_API(uint, blocker_shadow_set)
NODE_SOCKET_API(uint64_t, shadow_set_membership)
/* Set during device update. */
bool intersects_volume;

4
intern/cycles/scene/osl.cpp
View File

@ -1051,8 +1051,10 @@ void OSLCompiler::parameter(ShaderNode *node, const char *name)
case SocketType::CLOSURE:
case SocketType::NODE:
case SocketType::NODE_ARRAY:
case SocketType::UINT:
case SocketType::UINT64:
case SocketType::UNDEFINED:
case SocketType::UINT: {
case SocketType::NUM_TYPES: {
assert(0);
break;
}

6
intern/cycles/scene/scene.cpp
View File

@ -491,6 +491,12 @@ void Scene::update_kernel_features()
else if (geom->is_pointcloud()) {
kernel_features |= KERNEL_FEATURE_POINTCLOUD;
}
if (object->get_receiver_light_set()) {
kernel_features |= KERNEL_FEATURE_LIGHT_LINKING;
}
if (object->get_blocker_shadow_set()) {
kernel_features |= KERNEL_FEATURE_SHADOW_LINKING;
}
}
foreach (Light *light, lights) {

2
intern/cycles/scene/shader_nodes.h
View File

@ -443,7 +443,7 @@ class ConvertNode : public ShaderNode {
};
ustring value_string;
static const int MAX_TYPE = 12;
static const int MAX_TYPE = 13;
static bool register_types();
static Node *create(const NodeType *type);
static const NodeType *node_types[MAX_TYPE][MAX_TYPE];

2
intern/cycles/scene/stats.cpp
View File

@ -238,6 +238,7 @@ void RenderStats::collect_profiling(Scene *scene, Profiler &prof)
kernel.add_entry("Intersect Shadow", prof.get_event(PROFILING_INTERSECT_SHADOW));
kernel.add_entry("Intersect Subsurface", prof.get_event(PROFILING_INTERSECT_SUBSURFACE));
kernel.add_entry("Intersect Volume Stack", prof.get_event(PROFILING_INTERSECT_VOLUME_STACK));
kernel.add_entry("Intersect Blocked Light", prof.get_event(PROFILING_INTERSECT_DEDICATED_LIGHT));
NamedNestedSampleStats &surface = kernel.add_entry("Shade Surface", 0);
surface.add_entry("Setup", prof.get_event(PROFILING_SHADE_SURFACE_SETUP));
@ -257,6 +258,7 @@ void RenderStats::collect_profiling(Scene *scene, Profiler &prof)
shadow.add_entry("Setup", prof.get_event(PROFILING_SHADE_SHADOW_SETUP));
shadow.add_entry("Surface", prof.get_event(PROFILING_SHADE_SHADOW_SURFACE));
shadow.add_entry("Volume", prof.get_event(PROFILING_SHADE_SHADOW_VOLUME));
shadow.add_entry("Blocked Light", prof.get_event(PROFILING_SHADE_DEDICATED_LIGHT));
NamedNestedSampleStats &light = kernel.add_entry("Shade Light", 0);
light.add_entry("Setup", prof.get_event(PROFILING_SHADE_LIGHT_SETUP));

2
intern/cycles/util/profiling.h
View File

@ -20,6 +20,7 @@ enum ProfilingEvent : uint32_t {
PROFILING_INTERSECT_SUBSURFACE,
PROFILING_INTERSECT_SHADOW,
PROFILING_INTERSECT_VOLUME_STACK,
PROFILING_INTERSECT_DEDICATED_LIGHT,
PROFILING_SHADE_SURFACE_SETUP,
PROFILING_SHADE_SURFACE_EVAL,
@ -27,6 +28,7 @@ enum ProfilingEvent : uint32_t {
PROFILING_SHADE_SURFACE_INDIRECT_LIGHT,
PROFILING_SHADE_SURFACE_AO,
PROFILING_SHADE_SURFACE_PASSES,
PROFILING_SHADE_DEDICATED_LIGHT,
PROFILING_SHADE_VOLUME_SETUP,
PROFILING_SHADE_VOLUME_INTEGRATE,

2
intern/ghost/intern/GHOST_ContextVK.cc
View File

@ -824,7 +824,7 @@ static bool surfaceFormatSupported(const VkSurfaceFormatKHR &surface_format)
/**
* Select the surface format that we will use.
*
* We will select any 8bit unorm surface.
* We will select any 8bit UNORM surface.
*/
static bool selectSurfaceFormat(const VkPhysicalDevice physical_device,
const VkSurfaceKHR surface,

13
release/lts/lts_download.py
View File

@ -26,10 +26,11 @@ class Version:
def get_download_file_names(version: Version):
yield f"blender-{version}-linux-x64.tar.xz"
yield f"blender-{version}-macos-x64.dmg"
yield f"blender-{version}-windows-x64.msi"
yield f"blender-{version}-windows-x64.zip"
yield (f"blender-{version}-linux-x64.tar.xz", "Linux")
yield (f"blender-{version}-macos-x64.dmg", "macOS - Intel")
yield (f"blender-{version}-macos-arm64.dmg", "macOS - Apple Silicon")
yield (f"blender-{version}-windows-x64.msi", "Windows - Installer")
yield (f"blender-{version}-windows-x64.zip", "Windows - Portable (.zip)")
def get_download_url(version: Version, file_name: str) -> str:
@ -49,9 +50,9 @@ def generate_html(version: Version) -> str:
lines.append(f"Released on {today.strftime(DATE_FORMAT)}.")
lines.append("")
lines.append("<ul>")
for file_name in get_download_file_names(version):
for file_name, display_name in get_download_file_names(version):
download_url = get_download_url(version, file_name)
lines.append(f" <li><a href=\"{download_url}\">{file_name}</a></li>")
lines.append(f" <li><a href=\"{download_url}\">{display_name}</a></li>")
lines.append("</ul>")
return "\n".join(lines)

2
scripts/startup/bl_ui/space_time.py
View File

@ -257,7 +257,7 @@ class TIME_PT_playback(TimelinePanelButtons, Panel):
col.prop(screen, "use_play_clip_editors", text="Movie Clip Editor")
col.prop(screen, "use_play_node_editors", text="Node Editors")
col.prop(screen, "use_play_sequence_editors", text="Video Sequencer")
col.prop(screen, "use_play_spreadsheet_editors")
col.prop(screen, "use_play_spreadsheet_editors", text="Spreadsheet")
col = layout.column(heading="Show")
col.prop(scene, "show_subframe", text="Subframes")

4
scripts/startup/bl_ui/space_view3d.py
View File

@ -3046,6 +3046,10 @@ class VIEW3D_MT_make_links(Menu):
layout.operator("object.data_transfer")
layout.operator("object.datalayout_transfer")
layout.separator()
layout.operator_menu_enum("object.light_linking_receivers_link", "link_state")
layout.operator_menu_enum("object.light_linking_blockers_link", "link_state")
class VIEW3D_MT_brush_paint_modes(Menu):
bl_label = "Enabled Modes"

2
source/blender/blenfont/BLF_api.h
View File

@ -14,7 +14,7 @@
extern "C" {
#endif
/* Name of subfolder inside BLENDER_DATAFILES that contains font files. */
/* Name of sub-directory inside #BLENDER_DATAFILES that contains font files. */
#define BLF_DATAFILES_FONTS_DIR "fonts"
/* File name of the default variable-width font. */

2
source/blender/blenkernel/BKE_blender_version.h
View File

@ -18,7 +18,7 @@ extern "C" {
/* Blender major and minor version. */
#define BLENDER_VERSION 400
/* Blender patch version for bugfix releases. */
/* Blender patch version for bug-fix releases. */
#define BLENDER_VERSION_PATCH 0
/** Blender release cycle stage: alpha/beta/rc/release. */
#define BLENDER_VERSION_CYCLE alpha

1
source/blender/blenkernel/BKE_collection.h
View File

@ -26,6 +26,7 @@ struct BlendExpander;
struct BlendLibReader;
struct BlendWriter;
struct Collection;
struct Depsgraph;
struct ID;
struct Library;
struct Main;

2
source/blender/blenkernel/BKE_idprop.h
View File

@ -86,8 +86,6 @@ struct IDProperty *IDP_NewString(const char *st, const char *name) ATTR_WARN_UNU
*/
void IDP_AssignStringMaxSize(struct IDProperty *prop, const char *st, int maxncpy) ATTR_NONNULL();
void IDP_AssignString(struct IDProperty *prop, const char *st) ATTR_NONNULL();
void IDP_ConcatStringC(struct IDProperty *prop, const char *st) ATTR_NONNULL();
void IDP_ConcatString(struct IDProperty *str1, struct IDProperty *append) ATTR_NONNULL();
void IDP_FreeString(struct IDProperty *prop) ATTR_NONNULL();
/*-------- ID Type -------*/

4
source/blender/blenkernel/BKE_layer.h
View File

@ -16,10 +16,6 @@
extern "C" {
#endif
#define TODO_LAYER_OVERRIDE /* CollectionOverride */
#define TODO_LAYER_OPERATORS /* collection mamanger and property panel operators */
#define TODO_LAYER /* generic todo */
struct Base;
struct BlendDataReader;
struct BlendLibReader;

106
source/blender/blenkernel/BKE_light_linking.h Normal file
View File

@ -0,0 +1,106 @@
/* SPDX-License-Identifier: GPL-2.0-or-later
* Copyright 2001-2023 Blender Foundation */
#pragma once
/** \file
* \ingroup bke
*
* API to manage light linking.
*/
#include "DNA_collection_types.h" /* eCollectionLightLinkingState */
#ifdef __cplusplus
extern "C" {
#endif
struct ID;
struct Main;
struct Object;
struct Collection;
struct ReportList;
struct Scene;
struct ViewLayer;
typedef enum LightLinkingType {
LIGHT_LINKING_RECEIVER,
LIGHT_LINKING_BLOCKER,
} LightLinkingType;
/* Get a collection of the given light linking type of the given object. */
struct Collection *BKE_light_linking_collection_get(const struct Object *object,
LightLinkingType link_type);
/* Create new collection and assign it as a light or shadow linking collection (denoted by the
* link_type) of the given object.
*
* The collection is created outside of the view layer collections.
* If the object already has light linking collection set up it is unreferenced from the object.
*
* Returns the newly created collection. */
struct Collection *BKE_light_linking_collection_new(struct Main *bmain,
struct Object *object,
LightLinkingType link_type);
/* Assign given light or shadow linking collection (denoted by the link_type) to the given object.
* Maintains user counters of the collection: old collection is decreased the user counter, the new
* one is increased after this call.
* The new_collection is allowed to be null pointer.
*
* The assign_only variant takes care of (re)assigning the collection and maintaining the user
* counter, but not the dependency graph tagging for update. */
void BKE_light_linking_collection_assign_only(struct Object *object,
struct Collection *new_collection,
LightLinkingType link_type);
void BKE_light_linking_collection_assign(struct Main *bmain,
struct Object *object,
struct Collection *new_collection,
LightLinkingType link_type);
/* Add receiver to the given light linking collection.
* The ID is expected to either be collection or an object. Passing other types of IDs has no
* effect */
void BKE_light_linking_add_receiver_to_collection(struct Main *bmain,
struct Collection *collection,
struct ID *receiver,
const eCollectionLightLinkingState link_state);
/* Remove the given ID from the light or shadow linking collection of the given object.
*
* The collection is expected to be either receiver_collection or blocker_collection from an
* emitter object.
*
* The ID is expected to either be collection or an object. If other ID type is passed to the
* function an error is reported and false is returned.
*
* Returns true if the ID was unlinked from the receiver collection, false otherwise. The unlinking
* will be unsuccessful if, for example, the receiver collection is a linked data-block.
*
* The optional reports argument is used to provide human-readable details about why unlinking was
* not successful. If it is nullptr then the report is printed to the console. */
bool BKE_light_linking_unlink_id_from_collection(struct Main *bmain,
struct Collection *collection,
struct ID *id,
struct ReportList *reports);
/* Link receiver object to the given emitter.
*
* If the emitter already has light linking collection specified the object is added to that
* collection. Otherwise, first a new collection is created and assigned, and the receiver is added
* to it. */
void BKE_light_linking_link_receiver_to_emitter(struct Main *bmain,
struct Object *emitter,
struct Object *receiver,
LightLinkingType link_type,
eCollectionLightLinkingState link_state);
/* Select all objects which are linked to the given emitter via the given light link type. */
void BKE_light_linking_select_receivers_of_emitter(struct Scene *scene,
struct ViewLayer *view_layer,
struct Object *emitter,
LightLinkingType link_type);
#ifdef __cplusplus
}
#endif

97
source/blender/blenkernel/BKE_mesh_mapping.h
View File

@ -119,27 +119,6 @@ void BKE_mesh_uv_vert_map_free(UvVertMap *vmap);
#ifdef __cplusplus
/**
* Generates a map where the key is the vertex and the value
* is a list of polys that use that vertex as a corner.
* The lists are allocated from one memory pool.
*/
void BKE_mesh_vert_poly_map_create(MeshElemMap **r_map,
int **r_mem,
blender::OffsetIndices<int> polys,
const int *corner_verts,
int totvert);
/**
* Generates a map where the key is the vertex and the value
* is a list of loops that use that vertex as a corner.
* The lists are allocated from one memory pool.
*/
void BKE_mesh_vert_loop_map_create(MeshElemMap **r_map,
int **r_mem,
blender::OffsetIndices<int> polys,
const int *corner_verts,
int totvert);
/**
* Generates a map where the key is the edge and the value
* is a list of looptris that use that edge.
@ -152,42 +131,6 @@ void BKE_mesh_vert_looptri_map_create(MeshElemMap **r_map,
int totlooptri,
const int *corner_verts,
int totloop);
/**
* Generates a map where the key is the vertex and the value
* is a list of edges that use that vertex as an endpoint.
* The lists are allocated from one memory pool.
*/
void BKE_mesh_vert_edge_map_create(
MeshElemMap **r_map, int **r_mem, const blender::int2 *edges, int totvert, int totedge);
/**
* A version of #BKE_mesh_vert_edge_map_create that references connected vertices directly
* (not their edges).
*/
void BKE_mesh_vert_edge_vert_map_create(
MeshElemMap **r_map, int **r_mem, const blender::int2 *edges, int totvert, int totedge);
/**
* Generates a map where the key is the edge and the value is a list of loops that use that edge.
* Loops indices of a same poly are contiguous and in winding order.
* The lists are allocated from one memory pool.
*/
void BKE_mesh_edge_loop_map_create(MeshElemMap **r_map,
int **r_mem,
int totedge,
blender::OffsetIndices<int> polys,
const int *corner_edges,
int totloop);
/**
* Generates a map where the key is the edge and the value
* is a list of polygons that use that edge.
* The lists are allocated from one memory pool.
*/
void BKE_mesh_edge_poly_map_create(MeshElemMap **r_map,
int **r_mem,
int totedge,
blender::OffsetIndices<int> polys,
const int *corner_edges,
int totloop);
/**
* This function creates a map so the source-data (vert/edge/loop/poly)
* can loop over the destination data (using the destination arrays origindex).
@ -346,20 +289,36 @@ int *BKE_mesh_calc_smoothgroups(int totedge,
#ifdef __cplusplus
namespace blender::bke::mesh_topology {
namespace blender::bke::mesh {
Array<int> build_loop_to_poly_map(OffsetIndices<int> polys);
Array<Vector<int>> build_vert_to_edge_map(Span<int2> edges, int verts_num);
Array<Vector<int>> build_vert_to_poly_map(OffsetIndices<int> polys,
Span<int> corner_verts,
int verts_num);
Array<Vector<int>> build_vert_to_loop_map(Span<int> corner_verts, int verts_num);
Array<Vector<int>> build_edge_to_loop_map(Span<int> corner_edges, int edges_num);
Array<Vector<int, 2>> build_edge_to_poly_map(OffsetIndices<int> polys,
Span<int> corner_edges,
int edges_num);
Vector<Vector<int>> build_edge_to_loop_map_resizable(Span<int> corner_edges, int edges_num);
GroupedSpan<int> build_vert_to_edge_map(Span<int2> edges,
int verts_num,
Array<int> &r_offsets,
Array<int> &r_indices);
} // namespace blender::bke::mesh_topology
GroupedSpan<int> build_vert_to_poly_map(OffsetIndices<int> polys,
Span<int> corner_verts,
int verts_num,
Array<int> &r_offsets,
Array<int> &r_indices);
GroupedSpan<int> build_vert_to_loop_map(Span<int> corner_verts,
int verts_num,
Array<int> &r_offsets,
Array<int> &r_indices);
GroupedSpan<int> build_edge_to_loop_map(Span<int> corner_edges,
int edges_num,
Array<int> &r_offsets,
Array<int> &r_indices);
GroupedSpan<int> build_edge_to_poly_map(OffsetIndices<int> polys,
Span<int> corner_edges,
int edges_num,
Array<int> &r_offsets,
Array<int> &r_indices);
} // namespace blender::bke::mesh
#endif

3
source/blender/blenkernel/BKE_object.h
View File

@ -455,6 +455,9 @@ void BKE_object_eval_uber_transform(struct Depsgraph *depsgraph, struct Object *
void BKE_object_eval_uber_data(struct Depsgraph *depsgraph,
struct Scene *scene,
struct Object *ob);
void BKE_object_eval_light_linking(struct Depsgraph *depsgraph, struct Object *object);
/**
* Assign #Object.data after modifier stack evaluation.
*/

35
source/blender/blenkernel/BKE_paint.h
View File

@ -9,6 +9,10 @@
#include "BLI_bitmap.h"
#include "BLI_compiler_compat.h"
#ifdef __cplusplus
# include "BLI_array.hh"
# include "BLI_offset_indices.hh"
#endif
#include "BLI_utildefines.h"
#include "DNA_brush_enums.h"
@ -40,7 +44,6 @@ struct ListBase;
struct MLoopTri;
struct Main;
struct Mesh;
struct MeshElemMap;
struct Object;
struct PBVH;
struct Paint;
@ -273,12 +276,17 @@ void BKE_paint_blend_read_lib(struct BlendLibReader *reader,
#define SCULPT_FACE_SET_NONE 0
/** Used for both vertex color and weight paint. */
#ifdef __cplusplus
struct SculptVertexPaintGeomMap {
int *vert_map_mem;
struct MeshElemMap *vert_to_loop;
int *poly_map_mem;
struct MeshElemMap *vert_to_poly;
blender::Array<int> vert_to_loop_offsets;
blender::Array<int> vert_to_loop_indices;
blender::GroupedSpan<int> vert_to_loop;
blender::Array<int> vert_to_poly_offsets;
blender::Array<int> vert_to_poly_indices;
blender::GroupedSpan<int> vert_to_poly;
};
#endif
/** Pose Brush IK Chain. */
typedef struct SculptPoseIKChainSegment {
@ -599,16 +607,19 @@ typedef struct SculptSession {
/* Mesh connectivity maps. */
/* Vertices to adjacent polys. */
struct MeshElemMap *pmap;
int *pmap_mem;
blender::Array<int> vert_to_poly_offsets;
blender::Array<int> vert_to_poly_indices;
blender::GroupedSpan<int> pmap;
/* Edges to adjacent polys. */
struct MeshElemMap *epmap;
int *epmap_mem;
blender::Array<int> edge_to_poly_offsets;
blender::Array<int> edge_to_poly_indices;
blender::GroupedSpan<int> epmap;
/* Vertices to adjacent edges. */
struct MeshElemMap *vemap;
int *vemap_mem;
blender::Array<int> vert_to_edge_offsets;
blender::Array<int> vert_to_edge_indices;
blender::GroupedSpan<int> vemap;
/* Mesh Face Sets */
/* Total number of polys of the base mesh. */
@ -706,7 +717,7 @@ typedef struct SculptSession {
float prev_pivot_rot[4];
float prev_pivot_scale[3];
union {
struct {
struct {
struct SculptVertexPaintGeomMap gmap;
} vpaint;

4
source/blender/blenkernel/BKE_pbvh.h
View File

@ -35,7 +35,6 @@ struct DMFlagMat;
struct IsectRayPrecalc;
struct MLoopTri;
struct Mesh;
struct MeshElemMap;
struct PBVH;
struct PBVHBatches;
struct PBVHNode;
@ -45,7 +44,6 @@ struct SubdivCCG;
struct TaskParallelSettings;
struct Image;
struct ImageUser;
struct MeshElemMap;
typedef struct PBVH PBVH;
typedef struct PBVHNode PBVHNode;
@ -808,7 +806,6 @@ void BKE_pbvh_is_drawing_set(PBVH *pbvh, bool val);
void BKE_pbvh_node_num_loops(PBVH *pbvh, PBVHNode *node, int *r_totloop);
void BKE_pbvh_update_active_vcol(PBVH *pbvh, const struct Mesh *mesh);
void BKE_pbvh_pmap_set(PBVH *pbvh, const struct MeshElemMap *pmap);
void BKE_pbvh_vertex_color_set(PBVH *pbvh, PBVHVertRef vertex, const float color[4]);
void BKE_pbvh_vertex_color_get(const PBVH *pbvh, PBVHVertRef vertex, float r_color[4]);
@ -818,6 +815,7 @@ bool BKE_pbvh_draw_cache_invalid(const PBVH *pbvh);
int BKE_pbvh_debug_draw_gen_get(PBVHNode *node);
#ifdef __cplusplus
void BKE_pbvh_pmap_set(PBVH *pbvh, blender::GroupedSpan<int> pmap);
}
namespace blender::bke::pbvh {

2
source/blender/blenkernel/BKE_simulation_state.hh
View File

@ -162,7 +162,7 @@ class ModifierSimulationCache {
CacheState cache_state_ = CacheState::Valid;
bool failed_finding_bake_ = false;
void try_discover_bake(StringRefNull meta_dir, StringRefNull bdata_dir);
void try_discover_bake(StringRefNull absolute_bake_dir);
bool has_state_at_frame(const SubFrame &frame) const;
bool has_states() const;

9
source/blender/blenkernel/BKE_simulation_state_serialize.hh
View File

@ -143,12 +143,11 @@ class DiskBDataWriter : public BDataWriter {
};
/**
* Get the directory that contains all baked simulation data for the given modifier. This is a
* parent directory of the two directories below.
* Get the directory that contains all baked simulation data for the given modifier.
*/
std::string get_bake_directory(const Main &bmain, const Object &object, const ModifierData &md);
std::string get_bdata_directory(const Main &bmain, const Object &object, const ModifierData &md);
std::string get_meta_directory(const Main &bmain, const Object &object, const ModifierData &md);
std::string get_default_modifier_bake_directory(const Main &bmain,
const Object &object,
const ModifierData &md);
/**
* Encode the simulation state in a #DictionaryValue which also contains references to external

4
source/blender/blenkernel/BKE_subdiv_mesh.hh
View File

@ -8,10 +8,10 @@
#pragma once
#include "BLI_math_vector_types.hh"
#include "BLI_offset_indices.hh"
#include "BLI_sys_types.h"
struct Mesh;
struct MeshElemMap;
struct Subdiv;
struct SubdivToMeshSettings {
@ -36,7 +36,7 @@ Mesh *BKE_subdiv_to_mesh(Subdiv *subdiv,
* interpolation will be done base on the edge vertices. */
void BKE_subdiv_mesh_interpolate_position_on_edge(const float (*coarse_positions)[3],
const blender::int2 *coarse_edges,
const MeshElemMap *vert_to_edge_map,
blender::GroupedSpan<int> vert_to_edge_map,
int coarse_edge_index,
bool is_simple,
float u,

2
source/blender/blenkernel/CMakeLists.txt
View File

@ -181,6 +181,7 @@ set(SRC
intern/lib_remap.c
intern/library.c
intern/light.c
intern/light_linking.cc
intern/lightprobe.cc
intern/linestyle.cc
intern/main.c
@ -414,6 +415,7 @@ set(SRC
BKE_lib_remap.h
BKE_library.h
BKE_light.h
BKE_light_linking.h
BKE_lightprobe.h
BKE_linestyle.h
BKE_main.h

4
source/blender/blenkernel/intern/anim_sys.c
View File

@ -3250,8 +3250,8 @@ static void animsys_create_action_track_strip(const AnimData *adt,
* and this setting doesn't work. */
r_action_strip->flag |= NLASTRIP_FLAG_USR_INFLUENCE;
/* Unless extendmode is Nothing (might be useful for flattening NLA evaluation), disable range.
* Extendmode Nothing and Hold will behave as normal. Hold Forward will behave just like Hold.
/* Unless `extendmode` is Nothing (might be useful for flattening NLA evaluation), disable range.
* Extend-mode Nothing and Hold will behave as normal. Hold Forward will behave just like Hold.
*/
if (r_action_strip->extendmode != NLASTRIP_EXTEND_NOTHING) {
r_action_strip->flag |= NLASTRIP_FLAG_NO_TIME_MAP;

4
source/blender/blenkernel/intern/armature.c
View File

@ -104,7 +104,7 @@ static void armature_copy_data(Main *UNUSED(bmain), ID *id_dst, const ID *id_src
BLI_duplicatelist(&armature_dst->bonebase, &armature_src->bonebase);
/* Duplicate the childrens' lists */
/* Duplicate the children's lists. */
bone_dst = armature_dst->bonebase.first;
for (bone_src = armature_src->bonebase.first; bone_src; bone_src = bone_src->next) {
bone_dst->parent = NULL;
@ -2493,7 +2493,7 @@ void BKE_pose_where_is_bone(struct Depsgraph *depsgraph,
/* pose_mat(b) = pose_mat(b-1) * yoffs(b-1) * d_root(b) * bone_mat(b) * chan_mat(b) */
BKE_armature_mat_bone_to_pose(pchan, pchan->chan_mat, pchan->pose_mat);
/* Only rootbones get the cyclic offset (unless user doesn't want that). */
/* Only root-bones get the cyclic offset (unless user doesn't want that). */
/* XXX That could be a problem for snapping and other "reverse transform" features... */
if (!pchan->parent) {
if ((pchan->bone->flag & BONE_NO_CYCLICOFFSET) == 0) {

4
source/blender/blenkernel/intern/collection.c
View File

@ -970,7 +970,7 @@ static bool collection_object_cyclic_check_internal(Object *object, Collection *
if (object->instance_collection) {
Collection *dup_collection = object->instance_collection;
if ((dup_collection->id.tag & LIB_TAG_DOIT) == 0) {
/* Cycle already exists in collections, let's prevent further crappyness */
/* Cycle already exists in collections, let's prevent further creepiness. */
return true;
}
/* flag the object to identify cyclic dependencies in further dupli collections */
@ -1307,8 +1307,8 @@ static Collection *collection_parent_editable_find_recursive(const ViewLayer *vi
static bool collection_object_add(
Main *bmain, Collection *collection, Object *ob, int flag, const bool add_us)
{
/* Cyclic dependency check. */
if (ob->instance_collection) {
/* Cyclic dependency check. */
if ((ob->instance_collection == collection) ||
collection_find_child_recursive(ob->instance_collection, collection))
{

6
source/blender/blenkernel/intern/constraint.c
View File

@ -4025,7 +4025,7 @@ static void transform_evaluate(bConstraint *con, bConstraintOb *cob, ListBase *t
mat4_to_size(dvec, ct->matrix);
if (is_negative_m4(ct->matrix)) {
/* Bugfix #27886: (this is a limitation that riggers will have to live with for now).
/* Bug-fix #27886: (this is a limitation that riggers will have to live with for now).
* We can't be sure which axis/axes are negative,
* though we know that something is negative.
* Assume we don't care about negativity of separate axes. */
@ -6565,8 +6565,8 @@ void BKE_constraint_blend_read_lib(BlendLibReader *reader, ID *id, ListBase *con
/* legacy fixes */
LISTBASE_FOREACH (bConstraint *, con, conlist) {
/* patch for error introduced by changing constraints (dunno how) */
/* if con->data type changes, dna cannot resolve the pointer! (ton) */
/* Patch for error introduced by changing constraints (don't know how). */
/* NOTE(@ton): If `con->data` type changes, DNA cannot resolve the pointer!. */
if (con->data == NULL) {
con->type = CONSTRAINT_TYPE_NULL;
}

2
source/blender/blenkernel/intern/dynamicpaint.cc
View File

@ -97,7 +97,7 @@ static int neighStraightY[8] = {0, 1, 0, -1, 1, 1, -1, -1};
/* subframe_updateObject() flags */
#define SUBFRAME_RECURSION 5
/* surface_getBrushFlags() return vals */
/* #surface_getBrushFlags() return values. */
#define BRUSH_USES_VELOCITY (1 << 0)
/* Brush mesh ray-cast status. */
#define HIT_VOLUME 1

5
source/blender/blenkernel/intern/gpencil_modifier_legacy.c
View File

@ -38,6 +38,7 @@
#include "BKE_material.h"
#include "BKE_modifier.h"
#include "BKE_object.h"
#include "BKE_screen.h"
#include "BKE_shrinkwrap.h"
#include "DEG_depsgraph.h"
@ -447,9 +448,7 @@ const GpencilModifierTypeInfo *BKE_gpencil_modifier_get_info(GpencilModifierType
void BKE_gpencil_modifierType_panel_id(GpencilModifierType type, char *r_idname)
{
const GpencilModifierTypeInfo *mti = BKE_gpencil_modifier_get_info(type);
strcpy(r_idname, GPENCIL_MODIFIER_TYPE_PANEL_PREFIX);
strcat(r_idname, mti->name);
BLI_string_join(r_idname, BKE_ST_MAXNAME, GPENCIL_MODIFIER_TYPE_PANEL_PREFIX, mti->name);
}
void BKE_gpencil_modifier_panel_expand(GpencilModifierData *md)

Some files were not shown because too many files have changed in this diff Show More