blender
/
blender
133
398
Fork 578
Code Issues 6.3k Pull Requests 745 Projects 19 Wiki Activity

WIP: Vulkan: Workbench #107886

Closed
Jeroen Bakker wants to merge 88 commits from Jeroen-Bakker:vulkan-draw-manager-workbench into main
pull from: Jeroen-Bakker:vulkan-draw-manager-workbench
merge into: blender:main

When changing the target branch, be careful to rebase the branch in your fork to match. See documentation.
Conversation 0 Commits 88 Files Changed 20 +13908 -5880
602 changed files with 13908 additions and 5880 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files
Showing only changes of commit 2ce6e73570 - Show all commits

12
CMakeLists.txt
View File

@ -159,6 +159,15 @@ endif()
get_blender_version()
if(WIN32)
add_definitions(
# This is the app ID used for file registration, given it's used from several modules
# there really is no nice way to get this information consistent without a global define.
-DBLENDER_WIN_APPID="blender.${BLENDER_VERSION_MAJOR}.${BLENDER_VERSION_MINOR}"
# This is the name that will be shown in the taskbar and OpenWith windows UI
-DBLENDER_WIN_APPID_FRIENDLY_NAME="Blender ${BLENDER_VERSION_MAJOR}.${BLENDER_VERSION_MINOR}"
)
endif()
# -----------------------------------------------------------------------------
# Declare Options
@ -445,6 +454,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()

15
build_files/utils/make_update.py
View File

@ -445,10 +445,17 @@ def external_scripts_update(args: argparse.Namespace,
# automatically and fails when the branch is available in multiple remotes.
if make_utils.git_local_branch_exists(args.git_command, submodule_branch):
call([args.git_command, "checkout", submodule_branch])
elif make_utils.git_remote_exist(args.git_command, "origin"):
call([args.git_command, "checkout", "-t", f"origin/{submodule_branch}"])
elif make_utils.git_remote_exist(args.git_command, "upstream"):
call([args.git_command, "checkout", "-t", f"upstream/{submodule_branch}"])
else:
if make_utils.git_remote_branch_exists(args.git_command, "origin", submodule_branch):
call([args.git_command, "checkout", "-t", f"origin/{submodule_branch}"])
elif make_utils.git_remote_exist(args.git_command, "upstream"):
# For the Github style of upstream workflow create a local branch from
# the upstream, but do not track it, so that we stick to the paradigm
# that no local branches are tracking upstream, preventing possible
# accidental commit to upstream.
call([args.git_command, "checkout", "-b", submodule_branch,
f"upstream/{submodule_branch}", "--no-track"])
# Don't use extra fetch since all remotes of interest have been already fetched
# some lines above.
skip_msg += work_tree_update(args, use_fetch=False)

9
build_files/utils/make_utils.py
View File

@ -60,11 +60,16 @@ def git_local_branch_exists(git_command: str, branch: str) -> bool:
)
def git_remote_branch_exists(git_command: str, remote: str, branch: str) -> bool:
return call([git_command, "rev-parse", "--verify", f"remotes/{remote}/{branch}"],
exit_on_error=False, silent=True) == 0
def git_branch_exists(git_command: str, branch: str) -> bool:
return (
git_local_branch_exists(git_command, branch) or
call([git_command, "rev-parse", "--verify", "remotes/upstream/" + branch], exit_on_error=False, silent=True) == 0 or
call([git_command, "rev-parse", "--verify", "remotes/origin/" + branch], exit_on_error=False, silent=True) == 0
git_remote_branch_exists(git_command, "upstream", branch) or
git_remote_branch_exists(git_command, "origin", branch)
)

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/blender/util.h
View File

@ -118,7 +118,7 @@ static inline BL::Mesh object_to_mesh(BL::BlendData & /*data*/,
if ((bool)mesh && subdivision_type == Mesh::SUBDIVISION_NONE) {
if (mesh.use_auto_smooth()) {
mesh.calc_normals_split();
mesh.split_faces(false);
mesh.split_faces();
}
mesh.calc_loop_triangles();

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:

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

@ -32,20 +32,14 @@ const char *kernel_type_as_string(MetalPipelineType pso_type)
return "";
}
bool kernel_has_intersection(DeviceKernel device_kernel)
{
return (device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST ||
device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW ||
device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE ||
device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK ||
device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE);
}
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:
@ -393,6 +387,11 @@ bool MetalKernelPipeline::should_use_binary_archive() const
}
}
if (use_metalrt && device_kernel_has_intersection(device_kernel)) {
/* Binary linked functions aren't supported in binary archives. */
return false;
}
if (pso_type == PSO_GENERIC) {
/* Archive the generic kernels. */
return true;
@ -581,7 +580,7 @@ void MetalKernelPipeline::compile()
[unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL]];
[unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL_PRIM]];
if (kernel_has_intersection(device_kernel)) {
if (device_kernel_has_intersection(device_kernel)) {
linked_functions = [[NSArray arrayWithArray:[unique_functions allObjects]]
sortedArrayUsingComparator:^NSComparisonResult(id<MTLFunction> f1, id<MTLFunction> f2) {
return [f1.label compare:f2.label];

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;
}
}

18
intern/cycles/hydra/field.cpp
View File

@ -26,12 +26,20 @@ class HdCyclesVolumeLoader : public VDBImageLoader {
HdCyclesVolumeLoader(const std::string &filePath, const std::string &gridName)
: VDBImageLoader(gridName)
{
/* Disable delay loading and file copying, this has poor performance on network drivers. */
/* Disable delay loading and file copying, this has poor performance on network drives. */
const bool delay_load = false;
openvdb::io::File file(filePath);
file.setCopyMaxBytes(0);
if (file.open(delay_load)) {
grid = file.readGrid(gridName);
try {
openvdb::io::File file(filePath);
file.setCopyMaxBytes(0);
if (file.open(delay_load)) {
grid = file.readGrid(gridName);
}
}
catch (const openvdb::IoError &e) {
VLOG_WARNING << "Error loading OpenVDB file: " << e.what();
}
catch (...) {
VLOG_WARNING << "Error loading OpenVDB file: Unknown error";
}
}
};

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

73
intern/cycles/kernel/camera/camera.h
View File

@ -12,19 +12,19 @@ CCL_NAMESPACE_BEGIN
/* Perspective Camera */
ccl_device float2 camera_sample_aperture(ccl_constant KernelCamera *cam, float u, float v)
ccl_device float2 camera_sample_aperture(ccl_constant KernelCamera *cam, const float2 rand)
{
float blades = cam->blades;
float2 bokeh;
if (blades == 0.0f) {
/* sample disk */
bokeh = concentric_sample_disk(u, v);
bokeh = concentric_sample_disk(rand);
}
else {
/* sample polygon */
float rotation = cam->bladesrotation;
bokeh = regular_polygon_sample(blades, rotation, u, v);
bokeh = regular_polygon_sample(blades, rotation, rand);
}
/* anamorphic lens bokeh */
@ -34,15 +34,13 @@ ccl_device float2 camera_sample_aperture(ccl_constant KernelCamera *cam, float u
}
ccl_device void camera_sample_perspective(KernelGlobals kg,
float raster_x,
float raster_y,
float lens_u,
float lens_v,
const float2 raster_xy,
const float2 rand_lens,
ccl_private Ray *ray)
{
/* create ray form raster position */
ProjectionTransform rastertocamera = kernel_data.cam.rastertocamera;
float3 raster = make_float3(raster_x, raster_y, 0.0f);
const float3 raster = float2_to_float3(raster_xy);
float3 Pcamera = transform_perspective(&rastertocamera, raster);
if (kernel_data.cam.have_perspective_motion) {
@ -71,14 +69,14 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
if (aperturesize > 0.0f) {
/* sample point on aperture */
float2 lensuv = camera_sample_aperture(&kernel_data.cam, lens_u, lens_v) * aperturesize;
float2 lens_uv = camera_sample_aperture(&kernel_data.cam, rand_lens) * aperturesize;
/* compute point on plane of focus */
float ft = kernel_data.cam.focaldistance / D.z;
float3 Pfocus = D * ft;
/* update ray for effect of lens */
P = make_float3(lensuv.x, lensuv.y, 0.0f);
P = float2_to_float3(lens_uv);
D = normalize(Pfocus - P);
}
@ -133,7 +131,7 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
float3 Px = Pnostereo;
float3 Dx = transform_perspective(&rastertocamera,
make_float3(raster_x + 1.0f, raster_y, 0.0f));
make_float3(raster.x + 1.0f, raster.y, 0.0f));
Dx = normalize(transform_direction(&cameratoworld, Dx));
spherical_stereo_transform(&kernel_data.cam, &Px, &Dx);
@ -144,7 +142,7 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
float3 Py = Pnostereo;
float3 Dy = transform_perspective(&rastertocamera,
make_float3(raster_x, raster_y + 1.0f, 0.0f));
make_float3(raster.x, raster.y + 1.0f, 0.0f));
Dy = normalize(transform_direction(&cameratoworld, Dy));
spherical_stereo_transform(&kernel_data.cam, &Py, &Dy);
@ -166,15 +164,13 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
/* Orthographic Camera */
ccl_device void camera_sample_orthographic(KernelGlobals kg,
float raster_x,
float raster_y,
float lens_u,
float lens_v,
const float2 raster_xy,
const float2 rand_lens,
ccl_private Ray *ray)
{
/* create ray form raster position */
ProjectionTransform rastertocamera = kernel_data.cam.rastertocamera;
float3 Pcamera = transform_perspective(&rastertocamera, make_float3(raster_x, raster_y, 0.0f));
float3 Pcamera = transform_perspective(&rastertocamera, float2_to_float3(raster_xy));
float3 P;
float3 D = make_float3(0.0f, 0.0f, 1.0f);
@ -184,15 +180,15 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg,
if (aperturesize > 0.0f) {
/* sample point on aperture */
float2 lensuv = camera_sample_aperture(&kernel_data.cam, lens_u, lens_v) * aperturesize;
float2 lens_uv = camera_sample_aperture(&kernel_data.cam, rand_lens) * aperturesize;
/* compute point on plane of focus */
float3 Pfocus = D * kernel_data.cam.focaldistance;
/* update ray for effect of lens */
float3 lensuvw = make_float3(lensuv.x, lensuv.y, 0.0f);
P = Pcamera + lensuvw;
D = normalize(Pfocus - lensuvw);
float3 lens_uvw = float2_to_float3(lens_uv);
P = Pcamera + lens_uvw;
D = normalize(Pfocus - lens_uvw);
}
else {
P = Pcamera;
@ -229,14 +225,12 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg,
ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
ccl_global const DecomposedTransform *cam_motion,
float raster_x,
float raster_y,
float lens_u,
float lens_v,
const float2 raster,
const float2 rand_lens,
ccl_private Ray *ray)
{
ProjectionTransform rastertocamera = cam->rastertocamera;
float3 Pcamera = transform_perspective(&rastertocamera, make_float3(raster_x, raster_y, 0.0f));
float3 Pcamera = transform_perspective(&rastertocamera, float2_to_float3(raster));
/* create ray form raster position */
float3 P = zero_float3();
@ -253,7 +247,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
if (aperturesize > 0.0f) {
/* sample point on aperture */
float2 lensuv = camera_sample_aperture(cam, lens_u, lens_v) * aperturesize;
float2 lens_uv = camera_sample_aperture(cam, rand_lens) * aperturesize;
/* compute point on plane of focus */
float3 Dfocus = normalize(D);
@ -265,7 +259,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
V = normalize(cross(Dfocus, U));
/* update ray for effect of lens */
P = U * lensuv.x + V * lensuv.y;
P = U * lens_uv.x + V * lens_uv.y;
D = normalize(Pfocus - P);
}
@ -302,7 +296,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
Pcenter = transform_point(&cameratoworld, Pcenter);
Dcenter = normalize(transform_direction(&cameratoworld, Dcenter));
float3 Px = transform_perspective(&rastertocamera, make_float3(raster_x + 1.0f, raster_y, 0.0f));
float3 Px = transform_perspective(&rastertocamera, make_float3(raster.x + 1.0f, raster.y, 0.0f));
float3 Dx = panorama_to_direction(cam, Px.x, Px.y);
if (use_stereo) {
spherical_stereo_transform(cam, &Px, &Dx);
@ -314,7 +308,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
dP.dx = Px - Pcenter;
dD.dx = Dx - Dcenter;
float3 Py = transform_perspective(&rastertocamera, make_float3(raster_x, raster_y + 1.0f, 0.0f));
float3 Py = transform_perspective(&rastertocamera, make_float3(raster.x, raster.y + 1.0f, 0.0f));
float3 Dy = panorama_to_direction(cam, Py.x, Py.y);
if (use_stereo) {
spherical_stereo_transform(cam, &Py, &Dy);
@ -341,17 +335,16 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
ccl_device_inline void camera_sample(KernelGlobals kg,
int x,
int y,
float filter_u,
float filter_v,
float lens_u,
float lens_v,
float time,
const float2 filter_uv,
const float time,
const float2 lens_uv,
ccl_private Ray *ray)
{
/* pixel filter */
int filter_table_offset = kernel_data.tables.filter_table_offset;
float raster_x = x + lookup_table_read(kg, filter_u, filter_table_offset, FILTER_TABLE_SIZE);
float raster_y = y + lookup_table_read(kg, filter_v, filter_table_offset, FILTER_TABLE_SIZE);
const float2 raster = make_float2(
x + lookup_table_read(kg, filter_uv.x, filter_table_offset, FILTER_TABLE_SIZE),
y + lookup_table_read(kg, filter_uv.y, filter_table_offset, FILTER_TABLE_SIZE));
/* motion blur */
if (kernel_data.cam.shuttertime == -1.0f) {
@ -393,14 +386,14 @@ ccl_device_inline void camera_sample(KernelGlobals kg,
/* sample */
if (kernel_data.cam.type == CAMERA_PERSPECTIVE) {
camera_sample_perspective(kg, raster_x, raster_y, lens_u, lens_v, ray);
camera_sample_perspective(kg, raster, lens_uv, ray);
}
else if (kernel_data.cam.type == CAMERA_ORTHOGRAPHIC) {
camera_sample_orthographic(kg, raster_x, raster_y, lens_u, lens_v, ray);
camera_sample_orthographic(kg, raster, lens_uv, ray);
}
else {
ccl_global const DecomposedTransform *cam_motion = kernel_data_array(camera_motion);
camera_sample_panorama(&kernel_data.cam, cam_motion, raster_x, raster_y, lens_u, lens_v, ray);
camera_sample_panorama(&kernel_data.cam, cam_motion, raster, lens_uv, ray);
}
}

75
intern/cycles/kernel/closure/bsdf.h
View File

@ -122,73 +122,53 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
*pdf = 0.f;
int label = LABEL_NONE;
const float3 Ng = (sd->type & PRIMITIVE_CURVE) ? sc->N : sd->Ng;
const float2 rand_xy = float3_to_float2(rand);
switch (sc->type) {
case CLOSURE_BSDF_DIFFUSE_ID:
label = bsdf_diffuse_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_diffuse_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
#if defined(__SVM__) || defined(__OSL__)
case CLOSURE_BSDF_OREN_NAYAR_ID:
label = bsdf_oren_nayar_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_oren_nayar_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
# ifdef __OSL__
case CLOSURE_BSDF_PHONG_RAMP_ID:
label = bsdf_phong_ramp_sample(
sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
label = bsdf_phong_ramp_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
*eta = 1.0f;
break;
case CLOSURE_BSDF_DIFFUSE_RAMP_ID:
label = bsdf_diffuse_ramp_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_diffuse_ramp_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
# endif
case CLOSURE_BSDF_TRANSLUCENT_ID:
label = bsdf_translucent_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_translucent_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
case CLOSURE_BSDF_TRANSPARENT_ID:
label = bsdf_transparent_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_transparent_sample(sc, Ng, sd->wi, eval, wo, pdf);
*sampled_roughness = zero_float2();
*eta = 1.0f;
break;
case CLOSURE_BSDF_REFLECTION_ID:
case CLOSURE_BSDF_REFRACTION_ID:
case CLOSURE_BSDF_SHARP_GLASS_ID:
label = bsdf_microfacet_sharp_sample(sc,
path_flag,
Ng,
sd->wi,
rand.x,
rand.y,
rand.z,
eval,
wo,
pdf,
sampled_roughness,
eta);
label = bsdf_microfacet_sharp_sample(
sc, path_flag, Ng, sd->wi, rand, eval, wo, pdf, sampled_roughness, eta);
break;
case CLOSURE_BSDF_MICROFACET_GGX_ID:
case CLOSURE_BSDF_MICROFACET_GGX_CLEARCOAT_ID:
case CLOSURE_BSDF_MICROFACET_GGX_REFRACTION_ID:
case CLOSURE_BSDF_MICROFACET_GGX_GLASS_ID:
label = bsdf_microfacet_ggx_sample(sc,
path_flag,
Ng,
sd->wi,
rand.x,
rand.y,
rand.z,
eval,
wo,
pdf,
sampled_roughness,
eta);
label = bsdf_microfacet_ggx_sample(
sc, path_flag, Ng, sd->wi, rand, eval, wo, pdf, sampled_roughness, eta);
break;
case CLOSURE_BSDF_MICROFACET_MULTI_GGX_ID:
label = bsdf_microfacet_multi_ggx_sample(kg,
@ -221,61 +201,50 @@ ccl_device_inline int bsdf_sample(KernelGlobals kg,
case CLOSURE_BSDF_MICROFACET_BECKMANN_ID:
case CLOSURE_BSDF_MICROFACET_BECKMANN_REFRACTION_ID:
case CLOSURE_BSDF_MICROFACET_BECKMANN_GLASS_ID:
label = bsdf_microfacet_beckmann_sample(sc,
path_flag,
Ng,
sd->wi,
rand.x,
rand.y,
rand.z,
eval,
wo,
pdf,
sampled_roughness,
eta);
label = bsdf_microfacet_beckmann_sample(
sc, path_flag, Ng, sd->wi, rand, eval, wo, pdf, sampled_roughness, eta);
break;
case CLOSURE_BSDF_ASHIKHMIN_SHIRLEY_ID:
label = bsdf_ashikhmin_shirley_sample(
sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
*eta = 1.0f;
break;
case CLOSURE_BSDF_ASHIKHMIN_VELVET_ID:
label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_ashikhmin_velvet_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
case CLOSURE_BSDF_DIFFUSE_TOON_ID:
label = bsdf_diffuse_toon_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_diffuse_toon_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
case CLOSURE_BSDF_GLOSSY_TOON_ID:
label = bsdf_glossy_toon_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_glossy_toon_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
// double check if this is valid
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
case CLOSURE_BSDF_HAIR_REFLECTION_ID:
label = bsdf_hair_reflection_sample(
sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
*eta = 1.0f;
break;
case CLOSURE_BSDF_HAIR_TRANSMISSION_ID:
label = bsdf_hair_transmission_sample(
sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf, sampled_roughness);
sc, Ng, sd->wi, rand_xy, eval, wo, pdf, sampled_roughness);
*eta = 1.0f;
break;
case CLOSURE_BSDF_HAIR_PRINCIPLED_ID:
label = bsdf_principled_hair_sample(
kg, sc, sd, rand.x, rand.y, rand.z, eval, wo, pdf, sampled_roughness, eta);
label = bsdf_principled_hair_sample(kg, sc, sd, rand, eval, wo, pdf, sampled_roughness, eta);
break;
case CLOSURE_BSDF_PRINCIPLED_DIFFUSE_ID:
label = bsdf_principled_diffuse_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_principled_diffuse_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;
case CLOSURE_BSDF_PRINCIPLED_SHEEN_ID:
label = bsdf_principled_sheen_sample(sc, Ng, sd->wi, rand.x, rand.y, eval, wo, pdf);
label = bsdf_principled_sheen_sample(sc, Ng, sd->wi, rand_xy, eval, wo, pdf);
*sampled_roughness = one_float2();
*eta = 1.0f;
break;

45
intern/cycles/kernel/closure/bsdf_ashikhmin_shirley.h
View File

@ -111,24 +111,19 @@ ccl_device_forceinline Spectrum bsdf_ashikhmin_shirley_eval(ccl_private const Sh
return make_spectrum(out);
}
ccl_device_inline void bsdf_ashikhmin_shirley_sample_first_quadrant(float n_x,
float n_y,
float randu,
float randv,
ccl_private float *phi,
ccl_private float *cos_theta)
ccl_device_inline void bsdf_ashikhmin_shirley_sample_first_quadrant(
float n_x, float n_y, const float2 rand, ccl_private float *phi, ccl_private float *cos_theta)
{
*phi = atanf(sqrtf((n_x + 1.0f) / (n_y + 1.0f)) * tanf(M_PI_2_F * randu));
*phi = atanf(sqrtf((n_x + 1.0f) / (n_y + 1.0f)) * tanf(M_PI_2_F * rand.x));
float cos_phi = cosf(*phi);
float sin_phi = sinf(*phi);
*cos_theta = powf(randv, 1.0f / (n_x * cos_phi * cos_phi + n_y * sin_phi * sin_phi + 1.0f));
*cos_theta = powf(rand.y, 1.0f / (n_x * cos_phi * cos_phi + n_y * sin_phi * sin_phi + 1.0f));
}
ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -162,32 +157,28 @@ ccl_device int bsdf_ashikhmin_shirley_sample(ccl_private const ShaderClosure *sc
float cos_theta;
if (n_x == n_y) {
/* isotropic sampling */
phi = M_2PI_F * randu;
cos_theta = powf(randv, 1.0f / (n_x + 1.0f));
phi = M_2PI_F * rand.x;
cos_theta = powf(rand.y, 1.0f / (n_x + 1.0f));
}
else {
/* anisotropic sampling */
if (randu < 0.25f) { /* first quadrant */
float remapped_randu = 4.0f * randu;
bsdf_ashikhmin_shirley_sample_first_quadrant(
n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
if (rand.x < 0.25f) { /* first quadrant */
rand.x *= 4.0f;
bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
}
else if (randu < 0.5f) { /* second quadrant */
float remapped_randu = 4.0f * (.5f - randu);
bsdf_ashikhmin_shirley_sample_first_quadrant(
n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
else if (rand.x < 0.5f) { /* second quadrant */
rand.x = 4.0f * (0.5f - rand.x);
bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
phi = M_PI_F - phi;
}
else if (randu < 0.75f) { /* third quadrant */
float remapped_randu = 4.0f * (randu - 0.5f);
bsdf_ashikhmin_shirley_sample_first_quadrant(
n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
else if (rand.x < 0.75f) { /* third quadrant */
rand.x = 4.0f * (rand.x - 0.5f);
bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
phi = M_PI_F + phi;
}
else { /* fourth quadrant */
float remapped_randu = 4.0f * (1.0f - randu);
bsdf_ashikhmin_shirley_sample_first_quadrant(
n_x, n_y, remapped_randu, randv, &phi, &cos_theta);
rand.x = 4.0f * (1.0f - rand.x);
bsdf_ashikhmin_shirley_sample_first_quadrant(n_x, n_y, rand, &phi, &cos_theta);
phi = 2.0f * M_PI_F - phi;
}
}

5
intern/cycles/kernel/closure/bsdf_ashikhmin_velvet.h
View File

@ -78,8 +78,7 @@ ccl_device Spectrum bsdf_ashikhmin_velvet_eval(ccl_private const ShaderClosure *
ccl_device int bsdf_ashikhmin_velvet_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
const float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -90,7 +89,7 @@ ccl_device int bsdf_ashikhmin_velvet_sample(ccl_private const ShaderClosure *sc,
// we are viewing the surface from above - send a ray out with uniform
// distribution over the hemisphere
sample_uniform_hemisphere(N, randu, randv, wo, pdf);
sample_uniform_hemisphere(N, rand, wo, pdf);
if (!(dot(Ng, *wo) > 0)) {
*pdf = 0.0f;

10
intern/cycles/kernel/closure/bsdf_diffuse.h
View File

@ -42,8 +42,7 @@ ccl_device Spectrum bsdf_diffuse_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_diffuse_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -52,7 +51,7 @@ ccl_device int bsdf_diffuse_sample(ccl_private const ShaderClosure *sc,
float3 N = bsdf->N;
// distribution over the hemisphere
sample_cos_hemisphere(N, randu, randv, wo, pdf);
sample_cos_hemisphere(N, rand, wo, pdf);
if (dot(Ng, *wo) > 0.0f) {
*eval = make_spectrum(*pdf);
@ -88,8 +87,7 @@ ccl_device Spectrum bsdf_translucent_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_translucent_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -99,7 +97,7 @@ ccl_device int bsdf_translucent_sample(ccl_private const ShaderClosure *sc,
// we are viewing the surface from the right side - send a ray out with cosine
// distribution over the hemisphere
sample_cos_hemisphere(-N, randu, randv, wo, pdf);
sample_cos_hemisphere(-N, rand, wo, pdf);
if (dot(Ng, *wo) < 0) {
*eval = make_spectrum(*pdf);
}

5
intern/cycles/kernel/closure/bsdf_diffuse_ramp.h
View File

@ -67,8 +67,7 @@ ccl_device Spectrum bsdf_diffuse_ramp_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_diffuse_ramp_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -77,7 +76,7 @@ ccl_device int bsdf_diffuse_ramp_sample(ccl_private const ShaderClosure *sc,
float3 N = bsdf->N;
// distribution over the hemisphere
sample_cos_hemisphere(N, randu, randv, wo, pdf);
sample_cos_hemisphere(N, rand, wo, pdf);
if (dot(Ng, *wo) > 0.0f) {
*eval = rgb_to_spectrum(bsdf_diffuse_ramp_get_color(bsdf->colors, *pdf * M_PI_F) * M_1_PI_F);

14
intern/cycles/kernel/closure/bsdf_hair.h
View File

@ -143,8 +143,7 @@ ccl_device Spectrum bsdf_hair_transmission_eval(ccl_private const ShaderClosure
ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -165,7 +164,7 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
float a_R = fast_atan2f(((M_PI_2_F + theta_r) * 0.5f - offset) * roughness1_inv, 1.0f);
float b_R = fast_atan2f(((-M_PI_2_F + theta_r) * 0.5f - offset) * roughness1_inv, 1.0f);
float t = roughness1 * tanf(randu * (a_R - b_R) + b_R);
float t = roughness1 * tanf(rand.x * (a_R - b_R) + b_R);
float theta_h = t + offset;
float theta_i = 2 * theta_h - theta_r;
@ -173,7 +172,7 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
float costheta_i, sintheta_i;
fast_sincosf(theta_i, &sintheta_i, &costheta_i);
float phi = 2 * safe_asinf(1 - 2 * randv) * roughness2;
float phi = 2 * safe_asinf(1 - 2 * rand.y) * roughness2;
float phi_pdf = fast_cosf(phi * 0.5f) * 0.25f / roughness2;
@ -196,8 +195,7 @@ ccl_device int bsdf_hair_reflection_sample(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -219,7 +217,7 @@ ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc
float b_TT = fast_atan2f(((-M_PI_2_F + theta_r) / 2 - offset) * roughness1_inv, 1.0f);
float c_TT = 2 * fast_atan2f(M_PI_2_F / roughness2, 1.0f);
float t = roughness1 * tanf(randu * (a_TT - b_TT) + b_TT);
float t = roughness1 * tanf(rand.x * (a_TT - b_TT) + b_TT);
float theta_h = t + offset;
float theta_i = 2 * theta_h - theta_r;
@ -227,7 +225,7 @@ ccl_device int bsdf_hair_transmission_sample(ccl_private const ShaderClosure *sc
float costheta_i, sintheta_i;
fast_sincosf(theta_i, &sintheta_i, &costheta_i);
float p = roughness2 * tanf(c_TT * (randv - 0.5f));
float p = roughness2 * tanf(c_TT * (rand.y - 0.5f));
float phi = p + M_PI_F;
float theta_pdf = roughness1 /
(2 * (t * t + roughness1 * roughness1) * (a_TT - b_TT) * costheta_i);

21
intern/cycles/kernel/closure/bsdf_hair_principled.h
View File

@ -336,9 +336,7 @@ ccl_device Spectrum bsdf_principled_hair_eval(KernelGlobals kg,
ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
ccl_private const ShaderClosure *sc,
ccl_private ShaderData *sd,
float randu,
float randv,
float randw,
float3 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -381,12 +379,12 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
int p = 0;
for (; p < 3; p++) {
if (randw < Ap_energy[p]) {
if (rand.z < Ap_energy[p]) {
break;
}
randw -= Ap_energy[p];
rand.z -= Ap_energy[p];
}
randw /= Ap_energy[p];
rand.z /= Ap_energy[p];
float v = bsdf->v;
if (p == 1) {
@ -396,9 +394,10 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
v *= 4.0f;
}
randw = max(randw, 1e-5f);
const float fac = 1.0f + v * logf(randw + (1.0f - randw) * expf(-2.0f / v));
float sin_theta_i = -fac * sin_theta_o + cos_from_sin(fac) * cosf(M_2PI_F * randv) * cos_theta_o;
rand.z = max(rand.z, 1e-5f);
const float fac = 1.0f + v * logf(rand.z + (1.0f - rand.z) * expf(-2.0f / v));
float sin_theta_i = -fac * sin_theta_o +
cos_from_sin(fac) * cosf(M_2PI_F * rand.y) * cos_theta_o;
float cos_theta_i = cos_from_sin(sin_theta_i);
float angles[6];
@ -410,10 +409,10 @@ ccl_device int bsdf_principled_hair_sample(KernelGlobals kg,
float phi;
if (p < 3) {
phi = delta_phi(p, gamma_o, gamma_t) + sample_trimmed_logistic(randu, bsdf->s);
phi = delta_phi(p, gamma_o, gamma_t) + sample_trimmed_logistic(rand.x, bsdf->s);
}
else {
phi = M_2PI_F * randu;
phi = M_2PI_F * rand.x;
}
const float phi_i = phi_o + phi;

46
intern/cycles/kernel/closure/bsdf_microfacet.h
View File

@ -76,8 +76,7 @@ static_assert(sizeof(ShaderClosure) >= sizeof(MicrofacetBsdf), "MicrofacetBsdf i
ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
const float alpha_x,
const float alpha_y,
const float randu,
const float randv)
const float2 rand)
{
/* 1. stretch wi */
float3 wi_ = make_float3(alpha_x * wi.x, alpha_y * wi.y, wi.z);
@ -90,8 +89,8 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
if (wi_.z >= 0.99999f) {
/* Special case (normal incidence). */
const float r = sqrtf(-logf(randu));
const float phi = M_2PI_F * randv;
const float r = sqrtf(-logf(rand.x));
const float phi = M_2PI_F * rand.y;
slope_x = r * cosf(phi);
slope_y = r * sinf(phi);
}
@ -125,8 +124,8 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
* solve y = 1 + b + K * (1 - b * b)
*/
const float K = tan_theta_i * SQRT_PI_INV;
const float y_approx = randu * (1.0f + erf_a + K * (1 - erf_a * erf_a));
const float y_exact = randu * (1.0f + erf_a + K * exp_a2);
const float y_approx = rand.x * (1.0f + erf_a + K * (1 - erf_a * erf_a));
const float y_exact = rand.x * (1.0f + erf_a + K * exp_a2);
float b = K > 0 ? (0.5f - sqrtf(K * (K - y_approx + 1.0f) + 0.25f)) / K : y_approx - 1.0f;
float inv_erf = fast_ierff(b);
@ -155,7 +154,7 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
}
slope_x = inv_erf;
slope_y = fast_ierff(2.0f * randv - 1.0f);
slope_y = fast_ierff(2.0f * rand.y - 1.0f);
}
/* 3. rotate */
@ -178,8 +177,7 @@ ccl_device_forceinline float3 microfacet_beckmann_sample_vndf(const float3 wi,
ccl_device_forceinline float3 microfacet_ggx_sample_vndf(const float3 wi,
const float alpha_x,
const float alpha_y,
const float randu,
const float randv)
const float2 rand)
{
/* Section 3.2: Transforming the view direction to the hemisphere configuration. */
float3 wi_ = normalize(make_float3(alpha_x * wi.x, alpha_y * wi.y, wi.z));
@ -198,7 +196,7 @@ ccl_device_forceinline float3 microfacet_ggx_sample_vndf(const float3 wi,
}
/* Section 4.2: Parameterization of the projected area. */
float2 t = concentric_sample_disk(randu, randv);
float2 t = concentric_sample_disk(rand);
t.y = mix(safe_sqrtf(1.0f - sqr(t.x)), t.y, 0.5f * (1.0f + wi_.z));
/* Section 4.3: Reprojection onto hemisphere. */
@ -467,9 +465,7 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
const int path_flag,
float3 Ng,
float3 wi,
float randu,
float randv,
float randw,
const float3 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -513,11 +509,11 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
* space before and after sampling. */
local_I = make_float3(dot(X, wi), dot(Y, wi), cos_NI);
if (m_type == MicrofacetType::GGX) {
local_H = microfacet_ggx_sample_vndf(local_I, alpha_x, alpha_y, randu, randv);
local_H = microfacet_ggx_sample_vndf(local_I, alpha_x, alpha_y, float3_to_float2(rand));
}
else {
/* m_type == MicrofacetType::BECKMANN */
local_H = microfacet_beckmann_sample_vndf(local_I, alpha_x, alpha_y, randu, randv);
local_H = microfacet_beckmann_sample_vndf(local_I, alpha_x, alpha_y, float3_to_float2(rand));
}
H = X * local_H.x + Y * local_H.y + N * local_H.z;
@ -545,7 +541,7 @@ ccl_device int bsdf_microfacet_sample(ccl_private const ShaderClosure *sc,
* excessive noise for reflection highlights. */
float reflect_pdf = (path_flag & PATH_RAY_CAMERA) ? clamp(fresnel, 0.125f, 0.875f) :
fresnel;
do_refract = (randw >= reflect_pdf);
do_refract = (rand.z >= reflect_pdf);
lobe_pdf = do_refract ? (1.0f - reflect_pdf) : reflect_pdf;
}
}
@ -764,9 +760,7 @@ ccl_device int bsdf_microfacet_ggx_sample(ccl_private const ShaderClosure *sc,
const int path_flag,
float3 Ng,
float3 wi,
float randu,
float randv,
float randw,
const float3 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -774,7 +768,7 @@ ccl_device int bsdf_microfacet_ggx_sample(ccl_private const ShaderClosure *sc,
ccl_private float *eta)
{
return bsdf_microfacet_sample<MicrofacetType::GGX>(
sc, path_flag, Ng, wi, randu, randv, randw, eval, wo, pdf, sampled_roughness, eta);
sc, path_flag, Ng, wi, rand, eval, wo, pdf, sampled_roughness, eta);
}
/* Beckmann microfacet with Smith shadow-masking from:
@ -833,9 +827,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(ccl_private const ShaderClosure *
const int path_flag,
float3 Ng,
float3 wi,
float randu,
float randv,
float randw,
const float3 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -843,7 +835,7 @@ ccl_device int bsdf_microfacet_beckmann_sample(ccl_private const ShaderClosure *
ccl_private float *eta)
{
return bsdf_microfacet_sample<MicrofacetType::BECKMANN>(
sc, path_flag, Ng, wi, randu, randv, randw, eval, wo, pdf, sampled_roughness, eta);
sc, path_flag, Ng, wi, rand, eval, wo, pdf, sampled_roughness, eta);
}
/* Specular interface, not really a microfacet model but close enough that sharing code makes
@ -889,9 +881,7 @@ ccl_device int bsdf_microfacet_sharp_sample(ccl_private const ShaderClosure *sc,
const int path_flag,
float3 Ng,
float3 wi,
float randu,
float randv,
float randw,
const float3 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -899,7 +889,7 @@ ccl_device int bsdf_microfacet_sharp_sample(ccl_private const ShaderClosure *sc,
ccl_private float *eta)
{
return bsdf_microfacet_sample<MicrofacetType::SHARP>(
sc, path_flag, Ng, wi, randu, randv, randw, eval, wo, pdf, sampled_roughness, eta);
sc, path_flag, Ng, wi, rand, eval, wo, pdf, sampled_roughness, eta);
}
CCL_NAMESPACE_END

5
intern/cycles/kernel/closure/bsdf_oren_nayar.h
View File

@ -66,14 +66,13 @@ ccl_device Spectrum bsdf_oren_nayar_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_oren_nayar_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
{
ccl_private const OrenNayarBsdf *bsdf = (ccl_private const OrenNayarBsdf *)sc;
sample_uniform_hemisphere(bsdf->N, randu, randv, wo, pdf);
sample_uniform_hemisphere(bsdf->N, rand, wo, pdf);
if (dot(Ng, *wo) > 0.0f) {
*eval = bsdf_oren_nayar_get_intensity(sc, bsdf->N, wi, *wo);

7
intern/cycles/kernel/closure/bsdf_phong_ramp.h
View File

@ -78,8 +78,7 @@ ccl_device_inline float phong_ramp_exponent_to_roughness(float exponent)
ccl_device int bsdf_phong_ramp_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
const float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf,
@ -96,8 +95,8 @@ ccl_device int bsdf_phong_ramp_sample(ccl_private const ShaderClosure *sc,
float3 R = (2 * cosNI) * bsdf->N - wi;
float3 T, B;
make_orthonormals(R, &T, &B);
float phi = M_2PI_F * randu;
float cosTheta = powf(randv, 1 / (m_exponent + 1));
float phi = M_2PI_F * rand.x;
float cosTheta = powf(rand.y, 1 / (m_exponent + 1));
float sinTheta2 = 1 - cosTheta * cosTheta;
float sinTheta = sinTheta2 > 0 ? sqrtf(sinTheta2) : 0;
*wo = (cosf(phi) * sinTheta) * T + (sinf(phi) * sinTheta) * B + (cosTheta)*R;

5
intern/cycles/kernel/closure/bsdf_principled_diffuse.h
View File

@ -132,8 +132,7 @@ ccl_device Spectrum bsdf_principled_diffuse_eval(ccl_private const ShaderClosure
ccl_device int bsdf_principled_diffuse_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -142,7 +141,7 @@ ccl_device int bsdf_principled_diffuse_sample(ccl_private const ShaderClosure *s
float3 N = bsdf->N;
sample_cos_hemisphere(N, randu, randv, wo, pdf);
sample_cos_hemisphere(N, rand, wo, pdf);
if (dot(Ng, *wo) > 0) {
*eval = bsdf_principled_diffuse_compute_brdf(bsdf, N, wi, *wo, pdf);

5
intern/cycles/kernel/closure/bsdf_principled_sheen.h
View File

@ -84,8 +84,7 @@ ccl_device Spectrum bsdf_principled_sheen_eval(ccl_private const ShaderClosure *
ccl_device int bsdf_principled_sheen_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -94,7 +93,7 @@ ccl_device int bsdf_principled_sheen_sample(ccl_private const ShaderClosure *sc,
float3 N = bsdf->N;
sample_cos_hemisphere(N, randu, randv, wo, pdf);
sample_cos_hemisphere(N, rand, wo, pdf);
if (dot(Ng, *wo) > 0) {
float3 H = normalize(wi + *wo);

14
intern/cycles/kernel/closure/bsdf_toon.h
View File

@ -79,8 +79,7 @@ ccl_device Spectrum bsdf_diffuse_toon_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -89,10 +88,10 @@ ccl_device int bsdf_diffuse_toon_sample(ccl_private const ShaderClosure *sc,
float max_angle = bsdf->size * M_PI_2_F;
float smooth = bsdf->smooth * M_PI_2_F;
float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
float angle = sample_angle * randu;
float angle = sample_angle * rand.x;
if (sample_angle > 0.0f) {
sample_uniform_cone(bsdf->N, sample_angle, randu, randv, wo, pdf);
sample_uniform_cone(bsdf->N, sample_angle, rand, wo, pdf);
if (dot(Ng, *wo) > 0.0f) {
*eval = make_spectrum(*pdf * bsdf_toon_get_intensity(max_angle, smooth, angle));
@ -152,8 +151,7 @@ ccl_device Spectrum bsdf_glossy_toon_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_glossy_toon_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -168,9 +166,9 @@ ccl_device int bsdf_glossy_toon_sample(ccl_private const ShaderClosure *sc,
float3 R = (2 * cosNI) * bsdf->N - wi;
float sample_angle = bsdf_toon_get_sample_angle(max_angle, smooth);
float angle = sample_angle * randu;
float angle = sample_angle * rand.x;
sample_uniform_cone(R, sample_angle, randu, randv, wo, pdf);
sample_uniform_cone(R, sample_angle, rand, wo, pdf);
if (dot(Ng, *wo) > 0.0f) {
float cosNO = dot(bsdf->N, *wo);

2
intern/cycles/kernel/closure/bsdf_transparent.h
View File

@ -71,8 +71,6 @@ ccl_device Spectrum bsdf_transparent_eval(ccl_private const ShaderClosure *sc,
ccl_device int bsdf_transparent_sample(ccl_private const ShaderClosure *sc,
float3 Ng,
float3 wi,
float randu,
float randv,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)

19
intern/cycles/kernel/closure/volume.h
View File

@ -67,21 +67,20 @@ ccl_device Spectrum volume_henyey_greenstein_eval_phase(ccl_private const Shader
return make_spectrum(*pdf);
}
ccl_device float3
henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_private float *pdf)
ccl_device float3 henyey_greenstrein_sample(float3 D, float g, float2 rand, ccl_private float *pdf)
{
/* match pdf for small g */
float cos_theta;
bool isotropic = fabsf(g) < 1e-3f;
if (isotropic) {
cos_theta = (1.0f - 2.0f * randu);
cos_theta = (1.0f - 2.0f * rand.x);
if (pdf) {
*pdf = M_1_PI_F * 0.25f;
}
}
else {
float k = (1.0f - g * g) / (1.0f - g + 2.0f * g * randu);
float k = (1.0f - g * g) / (1.0f - g + 2.0f * g * rand.x);
cos_theta = (1.0f + g * g - k * k) / (2.0f * g);
if (pdf) {
*pdf = single_peaked_henyey_greenstein(cos_theta, g);
@ -89,7 +88,7 @@ henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_priva
}
float sin_theta = sin_from_cos(cos_theta);
float phi = M_2PI_F * randv;
float phi = M_2PI_F * rand.y;
float3 dir = make_float3(sin_theta * cosf(phi), sin_theta * sinf(phi), cos_theta);
float3 T, B;
@ -101,8 +100,7 @@ henyey_greenstrein_sample(float3 D, float g, float randu, float randv, ccl_priva
ccl_device int volume_henyey_greenstein_sample(ccl_private const ShaderVolumeClosure *svc,
float3 wi,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
@ -110,7 +108,7 @@ ccl_device int volume_henyey_greenstein_sample(ccl_private const ShaderVolumeClo
float g = svc->g;
/* note that wi points towards the viewer and so is used negated */
*wo = henyey_greenstrein_sample(-wi, g, randu, randv, pdf);
*wo = henyey_greenstrein_sample(-wi, g, rand, pdf);
*eval = make_spectrum(*pdf); /* perfect importance sampling */
return LABEL_VOLUME_SCATTER;
@ -128,13 +126,12 @@ ccl_device Spectrum volume_phase_eval(ccl_private const ShaderData *sd,
ccl_device int volume_phase_sample(ccl_private const ShaderData *sd,
ccl_private const ShaderVolumeClosure *svc,
float randu,
float randv,
float2 rand,
ccl_private Spectrum *eval,
ccl_private float3 *wo,
ccl_private float *pdf)
{
return volume_henyey_greenstein_sample(svc, sd->wi, randu, randv, eval, wo, pdf);
return volume_henyey_greenstein_sample(svc, sd->wi, rand, eval, wo, pdf);
}
/* Volume sampling utilities. */

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
}

13
intern/cycles/kernel/integrator/init_from_camera.h
View File

@ -32,16 +32,11 @@ ccl_device_inline void integrate_camera_sample(KernelGlobals kg,
path_rng_3D(kg, rng_hash, sample, PRNG_LENS_TIME) :
zero_float3();
const float rand_time = rand_time_lens.x;
const float2 rand_lens = make_float2(rand_time_lens.y, rand_time_lens.z);
/* Generate camera ray. */
camera_sample(kg,
x,
y,
rand_filter.x,
rand_filter.y,
rand_time_lens.y,
rand_time_lens.z,
rand_time_lens.x,
ray);
camera_sample(kg, x, y, rand_filter, rand_time, rand_lens, ray);
}
/* Return false to indicate that this pixel is finished.

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

252
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>
@ -156,12 +254,11 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
if (!light_sample_from_position(kg,
rng_state,
rand_light.z,
rand_light.x,
rand_light.y,
rand_light,
sd->time,
sd->P,
sd->N,
light_link_receiver_nee(kg, sd),
sd->flag,
bounce,
path_flag,
@ -241,12 +338,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 +357,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 +379,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 +493,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);
@ -529,7 +574,7 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
float3 ao_D;
float ao_pdf;
sample_cos_hemisphere(ao_N, rand_bsdf.x, rand_bsdf.y, &ao_D, &ao_pdf);
sample_cos_hemisphere(ao_N, rand_bsdf, &ao_D, &ao_pdf);
bool skip_self = true;
@ -546,6 +591,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 +604,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 +633,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 +688,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 +702,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 +742,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 +755,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 +777,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(

58
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__
@ -707,13 +710,12 @@ ccl_device_forceinline bool integrate_volume_equiangular_sample_light(
LightSample ls ccl_optional_struct_init;
if (!light_sample_from_volume_segment(kg,
rand_light.z,
rand_light.x,
rand_light.y,
rand_light,
sd->time,
sd->P,
ray->D,
ray->tmax - ray->tmin,
light_link_receiver_nee(kg, sd),
bounce,
path_flag,
&ls))
@ -772,12 +774,11 @@ ccl_device_forceinline void integrate_volume_direct_light(
if (!light_sample_from_position(kg,
rng_state,
rand_light.z,
rand_light.x,
rand_light.y,
rand_light,
sd->time,
P,
zero_float3(),
light_link_receiver_nee(kg, sd),
SD_BSDF_HAS_TRANSMISSION,
bounce,
path_flag,
@ -831,10 +832,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 +889,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 +984,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 +1193,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. */

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

@ -180,7 +180,7 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
/* Sample diffuse surface scatter into the object. */
float3 D;
float pdf;
sample_cos_hemisphere(-N, rand_bsdf.x, rand_bsdf.y, &D, &pdf);
sample_cos_hemisphere(-N, rand_bsdf, &D, &pdf);
if (dot(-Ng, D) <= 0.0f) {
return false;
}
@ -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. */
@ -325,8 +326,7 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
ray.D = newD;
}
else {
float3 newD = henyey_greenstrein_sample(
ray.D, anisotropy, rand_scatter.x, rand_scatter.y, &hg_pdf);
float3 newD = henyey_greenstrein_sample(ray.D, anisotropy, rand_scatter, &hg_pdf);
cos_theta = dot(newD, N);
ray.D = newD;
}

5
intern/cycles/kernel/integrator/volume_shader.h
View File

@ -317,8 +317,7 @@ ccl_device int volume_shader_phase_guided_sample(KernelGlobals kg,
else {
/* Sample phase. */
*phase_pdf = 0.0f;
label = volume_phase_sample(
sd, svc, rand_phase.x, rand_phase.y, &eval, wo, unguided_phase_pdf);
label = volume_phase_sample(sd, svc, rand_phase, &eval, wo, unguided_phase_pdf);
if (*unguided_phase_pdf != 0.0f) {
bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);
@ -357,7 +356,7 @@ ccl_device int volume_shader_phase_sample(KernelGlobals kg,
Spectrum eval = zero_spectrum();
*pdf = 0.0f;
int label = volume_phase_sample(sd, svc, rand_phase.x, rand_phase.y, &eval, wo, pdf);
int label = volume_phase_sample(sd, svc, rand_phase, &eval, wo, pdf);
if (*pdf != 0.0f) {
bsdf_eval_init(phase_eval, CLOSURE_VOLUME_HENYEY_GREENSTEIN_ID, eval);

31
intern/cycles/kernel/light/area.h
View File

@ -19,8 +19,7 @@ ccl_device_inline float area_light_rect_sample(float3 P,
const float len_u,
const float3 axis_v,
const float len_v,
float randu,
float randv,
const float2 rand,
bool sample_coord)
{
/* In our name system we're using P for the center, which is o in the paper. */
@ -59,7 +58,7 @@ ccl_device_inline float area_light_rect_sample(float3 P,
if (sample_coord) {
/* Compute cu. */
float au = randu * S + k;
float au = rand.x * S + k;
float fu = (cosf(au) * b0 - b1) / sinf(au);
float cu = 1.0f / sqrtf(fu * fu + b0sq) * (fu > 0.0f ? 1.0f : -1.0f);
cu = clamp(cu, -1.0f, 1.0f);
@ -73,7 +72,7 @@ ccl_device_inline float area_light_rect_sample(float3 P,
float d = sqrtf(xu * xu + z0sq);
float h0 = y0 / sqrtf(d * d + y0sq);
float h1 = y1 / sqrtf(d * d + y1sq);
float hv = h0 + randv * (h1 - h0), hv2 = hv * hv;
float hv = h0 + rand.y * (h1 - h0), hv2 = hv * hv;
float yv = (hv2 < 1.0f - 1e-6f) ? (hv * d) / sqrtf(1.0f - hv2) : y1;
/* Transform (xu, yv, z0) to world coords. */
@ -233,8 +232,7 @@ ccl_device_inline bool area_light_eval(const ccl_global KernelLight *klight,
const float3 ray_P,
ccl_private float3 *light_P,
ccl_private LightSample *ccl_restrict ls,
float randu,
float randv,
const float2 rand,
bool sample_coord)
{
float3 axis_u = klight->area.axis_u;
@ -250,9 +248,8 @@ ccl_device_inline bool area_light_eval(const ccl_global KernelLight *klight,
if (in_volume_segment) {
light_P_new += sample_rectangle ?
rectangle_sample(
axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, randu, randv) :
ellipse_sample(axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, randu, randv);
rectangle_sample(axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, rand) :
ellipse_sample(axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, rand);
ls->pdf = invarea;
}
else {
@ -273,7 +270,7 @@ ccl_device_inline bool area_light_eval(const ccl_global KernelLight *klight,
if (sample_rectangle) {
ls->pdf = area_light_rect_sample(
ray_P, &light_P_new, axis_u, len_u, axis_v, len_v, randu, randv, sample_coord);
ray_P, &light_P_new, axis_u, len_u, axis_v, len_v, rand, sample_coord);
}
else {
if (klight->area.tan_half_spread == 0.0f) {
@ -281,8 +278,7 @@ ccl_device_inline bool area_light_eval(const ccl_global KernelLight *klight,
}
else {
if (sample_coord) {
light_P_new += ellipse_sample(
axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, randu, randv);
light_P_new += ellipse_sample(axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, rand);
}
ls->pdf = 4.0f * M_1_PI_F / (len_u * len_v);
}
@ -311,8 +307,7 @@ ccl_device_inline bool area_light_eval(const ccl_global KernelLight *klight,
template<bool in_volume_segment>
ccl_device_inline bool area_light_sample(const ccl_global KernelLight *klight,
const float randu,
const float randv,
const float2 rand,
const float3 P,
ccl_private LightSample *ls)
{
@ -325,7 +320,7 @@ ccl_device_inline bool area_light_sample(const ccl_global KernelLight *klight,
}
}
if (!area_light_eval<in_volume_segment>(klight, P, &ls->P, ls, randu, randv, true)) {
if (!area_light_eval<in_volume_segment>(klight, P, &ls->P, ls, rand, true)) {
return false;
}
@ -358,11 +353,11 @@ ccl_device_forceinline void area_light_update_position(const ccl_global KernelLi
{
if (klight->area.tan_half_spread == 0) {
/* Update position on the light to keep the direction fixed. */
area_light_eval<false>(klight, P, &ls->P, ls, 0, 0, true);
area_light_eval<false>(klight, P, &ls->P, ls, zero_float2(), true);
}
else {
ls->D = normalize_len(ls->P - P, &ls->t);
area_light_eval<false>(klight, P, &ls->P, ls, 0, 0, false);
area_light_eval<false>(klight, P, &ls->P, ls, zero_float2(), false);
/* Convert pdf to be in area measure. */
ls->pdf /= lamp_light_pdf(ls->Ng, -ls->D, ls->t);
}
@ -421,7 +416,7 @@ ccl_device_inline bool area_light_sample_from_intersection(
ls->Ng = klight->area.dir;
float3 light_P = klight->co;
return area_light_eval<false>(klight, ray_P, &light_P, ls, 0, 0, false);
return area_light_eval<false>(klight, ray_P, &light_P, ls, zero_float2(), false);
}
template<bool in_volume_segment>

73
intern/cycles/kernel/light/background.h
View File

@ -10,10 +10,7 @@ CCL_NAMESPACE_BEGIN
/* Background Light */
ccl_device float3 background_map_sample(KernelGlobals kg,
float randu,
float randv,
ccl_private float *pdf)
ccl_device float3 background_map_sample(KernelGlobals kg, float2 rand, ccl_private float *pdf)
{
/* for the following, the CDF values are actually a pair of floats, with the
* function value as X and the actual CDF as Y. The last entry's function
@ -30,7 +27,7 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
int step = count >> 1;
int middle = first + step;
if (kernel_data_fetch(light_background_marginal_cdf, middle).y < randv) {
if (kernel_data_fetch(light_background_marginal_cdf, middle).y < rand.y) {
first = middle + 1;
count -= step + 1;
}
@ -46,7 +43,7 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
float2 cdf_last_v = kernel_data_fetch(light_background_marginal_cdf, res_y);
/* importance-sampled V direction */
float dv = inverse_lerp(cdf_v.y, cdf_next_v.y, randv);
float dv = inverse_lerp(cdf_v.y, cdf_next_v.y, rand.y);
float v = (index_v + dv) / res_y;
/* This is basically std::lower_bound as used by PBRT. */
@ -57,7 +54,8 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
int middle = first + step;
if (kernel_data_fetch(light_background_conditional_cdf, index_v * cdf_width + middle).y <
randu) {
rand.x)
{
first = middle + 1;
count -= step + 1;
}
@ -76,7 +74,7 @@ ccl_device float3 background_map_sample(KernelGlobals kg,
index_v * cdf_width + res_x);
/* importance-sampled U direction */
float du = inverse_lerp(cdf_u.y, cdf_next_u.y, randu);
float du = inverse_lerp(cdf_u.y, cdf_next_u.y, rand.x);
float u = (index_u + du) / res_x;
/* compute pdf */
@ -198,7 +196,7 @@ ccl_device_inline float background_portal_pdf(
}
else {
portal_pdf += area_light_rect_sample(
P, &lightpos, axis_u, len_u, axis_v, len_v, 0.0f, 0.0f, false);
P, &lightpos, axis_u, len_u, axis_v, len_v, zero_float2(), false);
}
}
@ -223,16 +221,15 @@ ccl_device int background_num_possible_portals(KernelGlobals kg, float3 P)
ccl_device float3 background_portal_sample(KernelGlobals kg,
float3 P,
float randu,
float randv,
float2 rand,
int num_possible,
ccl_private int *sampled_portal,
ccl_private float *pdf)
{
/* Pick a portal, then re-normalize randv. */
randv *= num_possible;
int portal = (int)randv;
randv -= portal;
/* Pick a portal, then re-normalize rand.y. */
rand.y *= num_possible;
int portal = (int)rand.y;
rand.y -= portal;
/* TODO(sergey): Some smarter way of finding portal to sample
* is welcome.
@ -255,14 +252,13 @@ ccl_device float3 background_portal_sample(KernelGlobals kg,
float3 D;
if (is_round) {
lightpos += ellipse_sample(axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, randu, randv);
lightpos += ellipse_sample(axis_u * len_u * 0.5f, axis_v * len_v * 0.5f, rand);
float t;
D = normalize_len(lightpos - P, &t);
*pdf = fabsf(klight->area.invarea) * lamp_light_pdf(dir, -D, t);
}
else {
*pdf = area_light_rect_sample(
P, &lightpos, axis_u, len_u, axis_v, len_v, randu, randv, true);
*pdf = area_light_rect_sample(P, &lightpos, axis_u, len_u, axis_v, len_v, rand, true);
D = normalize(lightpos - P);
}
@ -278,14 +274,13 @@ ccl_device float3 background_portal_sample(KernelGlobals kg,
}
ccl_device_inline float3 background_sun_sample(KernelGlobals kg,
float randu,
float randv,
float2 rand,
ccl_private float *pdf)
{
float3 D;
const float3 N = float4_to_float3(kernel_data.background.sun);
const float angle = kernel_data.background.sun.w;
sample_uniform_cone(N, angle, randu, randv, &D, pdf);
sample_uniform_cone(N, angle, rand, &D, pdf);
return D;
}
@ -296,8 +291,10 @@ ccl_device_inline float background_sun_pdf(KernelGlobals kg, float3 D)
return pdf_uniform_cone(N, D, angle);
}
ccl_device_inline float3 background_light_sample(
KernelGlobals kg, float3 P, float randu, float randv, ccl_private float *pdf)
ccl_device_inline float3 background_light_sample(KernelGlobals kg,
float3 P,
float2 rand,
ccl_private float *pdf)
{
float portal_method_pdf = kernel_data.background.portal_weight;
float sun_method_pdf = kernel_data.background.sun_weight;
@ -316,7 +313,7 @@ ccl_device_inline float3 background_light_sample(
if (pdf_fac == 0.0f) {
/* Use uniform as a fallback if we can't use any strategy. */
*pdf = 1.0f / M_4PI_F;
return sample_uniform_sphere(randu, randv);
return sample_uniform_sphere(rand);
}
pdf_fac = 1.0f / pdf_fac;
@ -325,23 +322,23 @@ ccl_device_inline float3 background_light_sample(
map_method_pdf *= pdf_fac;
/* We have 100% in total and split it between the three categories.
* Therefore, we pick portals if randu is between 0 and portal_method_pdf,
* sun if randu is between portal_method_pdf and (portal_method_pdf + sun_method_pdf)
* and map if randu is between (portal_method_pdf + sun_method_pdf) and 1. */
* Therefore, we pick portals if rand.x is between 0 and portal_method_pdf,
* sun if rand.x is between portal_method_pdf and (portal_method_pdf + sun_method_pdf)
* and map if rand.x is between (portal_method_pdf + sun_method_pdf) and 1. */
float sun_method_cdf = portal_method_pdf + sun_method_pdf;
int method = 0;
float3 D;
if (randu < portal_method_pdf) {
if (rand.x < portal_method_pdf) {
method = 0;
/* Rescale randu. */
/* Rescale rand.x. */
if (portal_method_pdf != 1.0f) {
randu /= portal_method_pdf;
rand.x /= portal_method_pdf;
}
/* Sample a portal. */
int portal;
D = background_portal_sample(kg, P, randu, randv, num_portals, &portal, pdf);
D = background_portal_sample(kg, P, rand, num_portals, &portal, pdf);
if (num_portals > 1) {
/* Ignore the chosen portal, its pdf is already included. */
*pdf += background_portal_pdf(kg, P, D, portal, NULL);
@ -353,14 +350,14 @@ ccl_device_inline float3 background_light_sample(
}
*pdf *= portal_method_pdf;
}
else if (randu < sun_method_cdf) {
else if (rand.x < sun_method_cdf) {
method = 1;
/* Rescale randu. */
/* Rescale rand.x. */
if (sun_method_pdf != 1.0f) {
randu = (randu - portal_method_pdf) / sun_method_pdf;
rand.x = (rand.x - portal_method_pdf) / sun_method_pdf;
}
D = background_sun_sample(kg, randu, randv, pdf);
D = background_sun_sample(kg, rand, pdf);
/* Skip MIS if this is the only method. */
if (sun_method_pdf == 1.0f) {
@ -370,12 +367,12 @@ ccl_device_inline float3 background_light_sample(
}
else {
method = 2;
/* Rescale randu. */
/* Rescale rand.x. */
if (map_method_pdf != 1.0f) {
randu = (randu - sun_method_cdf) / map_method_pdf;
rand.x = (rand.x - sun_method_cdf) / map_method_pdf;
}
D = background_map_sample(kg, randu, randv, pdf);
D = background_map_sample(kg, rand, pdf);
/* Skip MIS if this is the only method. */
if (map_method_pdf == 1.0f) {

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

@ -28,24 +28,24 @@ typedef struct LightSample {
/* Utilities */
ccl_device_inline float3 ellipse_sample(float3 ru, float3 rv, float randu, float randv)
ccl_device_inline float3 ellipse_sample(float3 ru, float3 rv, float2 rand)
{
const float2 rand = concentric_sample_disk(randu, randv);
return ru * rand.x + rv * rand.y;
const float2 uv = concentric_sample_disk(rand);
return ru * uv.x + rv * uv.y;
}
ccl_device_inline float3 rectangle_sample(float3 ru, float3 rv, float randu, float randv)
ccl_device_inline float3 rectangle_sample(float3 ru, float3 rv, float2 rand)
{
return ru * (2.0f * randu - 1.0f) + rv * (2.0f * randv - 1.0f);
return ru * (2.0f * rand.x - 1.0f) + rv * (2.0f * rand.y - 1.0f);
}
ccl_device float3 disk_light_sample(float3 v, float randu, float randv)
ccl_device float3 disk_light_sample(float3 n, float2 rand)
{
float3 ru, rv;
make_orthonormals(v, &ru, &rv);
make_orthonormals(n, &ru, &rv);
return ellipse_sample(ru, rv, randu, randv);
return ellipse_sample(ru, rv, rand);
}
ccl_device float lamp_light_pdf(const float3 Ng, const float3 I, float t)
@ -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

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

@ -10,8 +10,7 @@
CCL_NAMESPACE_BEGIN
ccl_device_inline bool distant_light_sample(const ccl_global KernelLight *klight,
const float randu,
const float randv,
const float2 rand,
ccl_private LightSample *ls)
{
/* distant light */
@ -21,7 +20,7 @@ ccl_device_inline bool distant_light_sample(const ccl_global KernelLight *klight
float invarea = klight->distant.invarea;
if (radius > 0.0f) {
D = normalize(D + disk_light_sample(D, randu, randv) * radius);
D = normalize(D + disk_light_sample(D, rand) * radius);
}
ls->P = D;
@ -36,6 +35,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,

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

@ -11,7 +11,7 @@ CCL_NAMESPACE_BEGIN
/* Simple CDF based sampling over all lights in the scene, without taking into
* account shading position or normal. */
ccl_device int light_distribution_sample(KernelGlobals kg, const float randn)
ccl_device int light_distribution_sample(KernelGlobals kg, const float rand)
{
/* This is basically std::upper_bound as used by PBRT, to find a point light or
* triangle to emit from, proportional to area. a good improvement would be to
@ -19,13 +19,12 @@ ccl_device int light_distribution_sample(KernelGlobals kg, const float randn)
* arbitrary shaders. */
int first = 0;
int len = kernel_data.integrator.num_distribution + 1;
float r = randn;
do {
int half_len = len >> 1;
int middle = first + half_len;
if (r < kernel_data_fetch(light_distribution, middle).totarea) {
if (rand < kernel_data_fetch(light_distribution, middle).totarea) {
len = half_len;
}
else {
@ -43,20 +42,21 @@ ccl_device int light_distribution_sample(KernelGlobals kg, const float randn)
template<bool in_volume_segment>
ccl_device_noinline bool light_distribution_sample(KernelGlobals kg,
const float randn,
const float randu,
const float randv,
const float3 rand,
const float time,
const float3 P,
const int object_receiver,
const int bounce,
const uint32_t path_flag,
ccl_private LightSample *ls)
{
/* Sample light index from distribution. */
const int index = light_distribution_sample(kg, randn);
/* The first two dimensions of the Sobol sequence have better stratification. */
const int index = light_distribution_sample(kg, rand.z);
const float pdf_selection = kernel_data.integrator.distribution_pdf_lights;
const float2 rand_uv = float3_to_float2(rand);
return light_sample<in_volume_segment>(
kg, randu, randv, time, P, bounce, path_flag, index, 0, pdf_selection, ls);
kg, rand_uv, time, P, object_receiver, bounce, path_flag, index, 0, pdf_selection, ls);
}
ccl_device_inline float light_distribution_pdf_lamp(KernelGlobals kg)

264
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,13 +22,78 @@ 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>
ccl_device_inline bool light_sample(KernelGlobals kg,
const int lamp,
const float randu,
const float randv,
const float2 rand,
const float3 P,
const uint32_t path_flag,
ccl_private LightSample *ls)
@ -45,8 +111,8 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
ls->object = PRIM_NONE;
ls->prim = PRIM_NONE;
ls->lamp = lamp;
ls->u = randu;
ls->v = randv;
ls->u = rand.x;
ls->v = rand.y;
ls->group = lamp_lightgroup(kg, lamp);
if (in_volume_segment && (type == LIGHT_DISTANT || type == LIGHT_BACKGROUND)) {
@ -63,13 +129,13 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
}
if (type == LIGHT_DISTANT) {
if (!distant_light_sample(klight, randu, randv, ls)) {
if (!distant_light_sample(klight, rand, ls)) {
return false;
}
}
else if (type == LIGHT_BACKGROUND) {
/* infinite area light (e.g. light dome or env light) */
float3 D = -background_light_sample(kg, P, randu, randv, &ls->pdf);
float3 D = -background_light_sample(kg, P, rand, &ls->pdf);
ls->P = D;
ls->Ng = D;
@ -78,18 +144,18 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
ls->eval_fac = 1.0f;
}
else if (type == LIGHT_SPOT) {
if (!spot_light_sample<in_volume_segment>(klight, randu, randv, P, ls)) {
if (!spot_light_sample<in_volume_segment>(klight, rand, P, ls)) {
return false;
}
}
else if (type == LIGHT_POINT) {
if (!point_light_sample<in_volume_segment>(klight, randu, randv, P, ls)) {
if (!point_light_sample<in_volume_segment>(klight, rand, P, ls)) {
return false;
}
}
else {
/* area light */
if (!area_light_sample<in_volume_segment>(klight, randu, randv, P, ls)) {
if (!area_light_sample<in_volume_segment>(klight, rand, P, ls)) {
return false;
}
}
@ -101,10 +167,10 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
template<bool in_volume_segment>
ccl_device_noinline bool light_sample(KernelGlobals kg,
const float randu,
const float randv,
const float2 rand,
const float time,
const float3 P,
const int object_receiver,
const int bounce,
const uint32_t path_flag,
const int emitter_index,
@ -138,6 +204,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))
@ -146,17 +216,23 @@ ccl_device_noinline bool light_sample(KernelGlobals kg,
}
const int shader_flag = mesh_light.shader_flag;
if (!triangle_light_sample<in_volume_segment>(kg, prim, object, randu, randv, time, ls, P)) {
if (!triangle_light_sample<in_volume_segment>(kg, prim, object, rand, time, ls, P)) {
return false;
}
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, rand, P, path_flag, ls)) {
return false;
}
}
@ -167,15 +243,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 +278,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 +290,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 +334,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,

5
intern/cycles/kernel/light/point.h
View File

@ -9,8 +9,7 @@ CCL_NAMESPACE_BEGIN
template<bool in_volume_segment>
ccl_device_inline bool point_light_sample(const ccl_global KernelLight *klight,
const float randu,
const float randv,
const float2 rand,
const float3 P,
ccl_private LightSample *ls)
{
@ -21,7 +20,7 @@ ccl_device_inline bool point_light_sample(const ccl_global KernelLight *klight,
ls->P = center;
if (radius > 0.0f) {
ls->P += disk_light_sample(lightN, randu, randv) * radius;
ls->P += disk_light_sample(lightN, rand) * radius;
}
ls->pdf = klight->spot.invarea;

66
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. */
@ -317,13 +318,12 @@ ccl_device_inline float light_sample_mis_weight_nee(KernelGlobals kg,
* Uses either a flat distribution or light tree. */
ccl_device_inline bool light_sample_from_volume_segment(KernelGlobals kg,
const float randn,
const float randu,
const float randv,
const float3 rand,
const float time,
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)
@ -331,24 +331,23 @@ ccl_device_inline bool light_sample_from_volume_segment(KernelGlobals kg,
#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);
kg, rand, 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, rand, time, P, object_receiver, bounce, path_flag, ls);
}
}
ccl_device bool light_sample_from_position(KernelGlobals kg,
ccl_private const RNGState *rng_state,
const float randn,
const float randu,
const float randv,
const float3 rand,
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,
@ -357,46 +356,13 @@ 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);
kg, rand, 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, rand, time, P, object_receiver, bounce, path_flag, ls);
}
}
@ -440,7 +406,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 +431,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 +472,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

5
intern/cycles/kernel/light/spot.h
View File

@ -18,8 +18,7 @@ ccl_device float spot_light_attenuation(const ccl_global KernelSpotLight *spot,
template<bool in_volume_segment>
ccl_device_inline bool spot_light_sample(const ccl_global KernelLight *klight,
const float randu,
const float randv,
const float2 rand,
const float3 P,
ccl_private LightSample *ls)
{
@ -33,7 +32,7 @@ ccl_device_inline bool spot_light_sample(const ccl_global KernelLight *klight,
if (radius > 0.0f) {
/* disk light */
ls->P += disk_light_sample(lightN, randu, randv) * radius;
ls->P += disk_light_sample(lightN, rand) * radius;
}
const float invarea = klight->spot.invarea;

80
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,15 +656,27 @@ 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,
const float randu,
const float randv,
const float3 rand,
const float time,
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 +690,10 @@ 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);
/* The first two dimensions of the Sobol sequence have better stratification. */
float rand_selection = rand.z;
float3 local_P = P;
@ -690,21 +704,20 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
if (is_leaf(knode)) {
/* At a leaf node, we pick an emitter. */
selected_emitter = light_tree_cluster_select_emitter<in_volume_segment>(
kg, randn, local_P, N_or_D, t, has_transmission, &node_index, &pdf_selection);
kg, rand_selection, local_P, N_or_D, t, has_transmission, &node_index, &pdf_selection);
if (node_index < 0) {
break;
}
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;
@ -717,7 +730,8 @@ ccl_device_noinline bool light_tree_sample(KernelGlobals kg,
float discard;
float total_prob = left_prob;
node_index = left_index;
sample_resevoir(right_index, 1.0f - left_prob, node_index, discard, total_prob, randn);
sample_resevoir(
right_index, 1.0f - left_prob, node_index, discard, total_prob, rand_selection);
pdf_leaf *= (node_index == left_index) ? left_prob : (1.0f - left_prob);
}
@ -727,25 +741,39 @@ 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,
float3_to_float2(rand),
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;

13
intern/cycles/kernel/light/triangle.h
View File

@ -122,8 +122,7 @@ template<bool in_volume_segment>
ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
int prim,
int object,
float randu,
float randv,
const float2 rand,
float time,
ccl_private LightSample *ls,
const float3 P)
@ -191,14 +190,14 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
/* precompute a few things
* these could be re-used to take several samples
* as they are independent of randu/randv */
* as they are independent of `rand` */
const float cos_c = dot(A, B);
const float sin_alpha = fast_sinf(alpha);
const float product = sin_alpha * cos_c;
/* Select a random sub-area of the spherical triangle
* and calculate the third vertex C_ of that new triangle */
const float phi = randu * solid_angle - alpha;
const float phi = rand.x * solid_angle - alpha;
float s, t;
fast_sincosf(phi, &s, &t);
const float u = t - cos_alpha;
@ -217,7 +216,7 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
/* Finally, select a random point along the edge of the new triangle
* That point on the spherical triangle is the sampled ray direction */
const float z = 1.0f - randv * (1.0f - dot(C_, B));
const float z = 1.0f - rand.y * (1.0f - dot(C_, B));
ls->D = z * B + sin_from_cos(z) * safe_normalize(C_ - dot(C_, B) * B);
/* calculate intersection with the planar triangle */
@ -246,8 +245,8 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
/* compute random point in triangle. From Eric Heitz's "A Low-Distortion Map Between Triangle
* and Square" */
float u = randu;
float v = randv;
float u = rand.x;
float v = rand.y;
if (v > u) {
u *= 0.5f;
v -= u;

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

65
intern/cycles/kernel/sample/mapping.h
View File

@ -11,13 +11,13 @@
CCL_NAMESPACE_BEGIN
/* distribute uniform xy on [0,1] over unit disk [-1,1] */
ccl_device void to_unit_disk(ccl_private float *x, ccl_private float *y)
ccl_device void to_unit_disk(ccl_private float2 *rand)
{
float phi = M_2PI_F * (*x);
float r = sqrtf(*y);
float phi = M_2PI_F * rand->x;
float r = sqrtf(rand->y);
*x = r * cosf(phi);
*y = r * sinf(phi);
rand->x = r * cosf(phi);
rand->y = r * sinf(phi);
}
/* return an orthogonal tangent and bitangent given a normal and tangent that
@ -32,24 +32,29 @@ ccl_device void make_orthonormals_tangent(const float3 N,
}
/* sample direction with cosine weighted distributed in hemisphere */
ccl_device_inline void sample_cos_hemisphere(
const float3 N, float randu, float randv, ccl_private float3 *wo, ccl_private float *pdf)
ccl_device_inline void sample_cos_hemisphere(const float3 N,
float2 rand,
ccl_private float3 *wo,
ccl_private float *pdf)
{
to_unit_disk(&randu, &randv);
float costheta = sqrtf(max(1.0f - randu * randu - randv * randv, 0.0f));
to_unit_disk(&rand);
float costheta = safe_sqrtf(1.0f - len_squared(rand));
float3 T, B;
make_orthonormals(N, &T, &B);
*wo = randu * T + randv * B + costheta * N;
*wo = rand.x * T + rand.y * B + costheta * N;
*pdf = costheta * M_1_PI_F;
}
/* sample direction uniformly distributed in hemisphere */
ccl_device_inline void sample_uniform_hemisphere(
const float3 N, float randu, float randv, ccl_private float3 *wo, ccl_private float *pdf)
ccl_device_inline void sample_uniform_hemisphere(const float3 N,
const float2 rand,
ccl_private float3 *wo,
ccl_private float *pdf)
{
float z = randu;
float r = sqrtf(max(0.0f, 1.0f - z * z));
float phi = M_2PI_F * randv;
float z = rand.x;
float r = sin_from_cos(z);
float phi = M_2PI_F * rand.y;
float x = r * cosf(phi);
float y = r * sinf(phi);
@ -60,17 +65,13 @@ ccl_device_inline void sample_uniform_hemisphere(
}
/* sample direction uniformly distributed in cone */
ccl_device_inline void sample_uniform_cone(const float3 N,
float angle,
float randu,
float randv,
ccl_private float3 *wo,
ccl_private float *pdf)
ccl_device_inline void sample_uniform_cone(
const float3 N, float angle, const float2 rand, ccl_private float3 *wo, ccl_private float *pdf)
{
const float cosThetaMin = cosf(angle);
const float cosTheta = mix(cosThetaMin, 1.0f, randu);
const float cosTheta = mix(cosThetaMin, 1.0f, rand.x);
const float sinTheta = sin_from_cos(cosTheta);
const float phi = M_2PI_F * randv;
const float phi = M_2PI_F * rand.y;
const float x = sinTheta * cosf(phi);
const float y = sinTheta * sinf(phi);
const float z = cosTheta;
@ -92,11 +93,11 @@ ccl_device_inline float pdf_uniform_cone(const float3 N, float3 D, float angle)
}
/* sample uniform point on the surface of a sphere */
ccl_device float3 sample_uniform_sphere(float u1, float u2)
ccl_device float3 sample_uniform_sphere(const float2 rand)
{
float z = 1.0f - 2.0f * u1;
float r = sqrtf(fmaxf(0.0f, 1.0f - z * z));
float phi = M_2PI_F * u2;
float z = 1.0f - 2.0f * rand.x;
float r = sin_from_cos(z);
float phi = M_2PI_F * rand.y;
float x = r * cosf(phi);
float y = r * sinf(phi);
@ -105,11 +106,11 @@ ccl_device float3 sample_uniform_sphere(float u1, float u2)
/* distribute uniform xy on [0,1] over unit disk [-1,1], with concentric mapping
* to better preserve stratification for some RNG sequences */
ccl_device float2 concentric_sample_disk(float u1, float u2)
ccl_device float2 concentric_sample_disk(const float2 rand)
{
float phi, r;
float a = 2.0f * u1 - 1.0f;
float b = 2.0f * u2 - 1.0f;
float a = 2.0f * rand.x - 1.0f;
float b = 2.0f * rand.y - 1.0f;
if (a == 0.0f && b == 0.0f) {
return zero_float2();
@ -127,8 +128,10 @@ ccl_device float2 concentric_sample_disk(float u1, float u2)
}
/* sample point in unit polygon with given number of corners and rotation */
ccl_device float2 regular_polygon_sample(float corners, float rotation, float u, float v)
ccl_device float2 regular_polygon_sample(float corners, float rotation, const float2 rand)
{
float u = rand.x, v = rand.y;
/* sample corner number and reuse u */
float corner = floorf(u * corners);
u = u * corners - corner;

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

@ -52,7 +52,7 @@ ccl_device float svm_ao(
const float2 rand_disk = path_branched_rng_2D(
kg, &rng_state, sample, num_samples, PRNG_SURFACE_AO);
float2 d = concentric_sample_disk(rand_disk.x, rand_disk.y);
float2 d = concentric_sample_disk(rand_disk);
float3 D = make_float3(d.x, d.y, safe_sqrtf(1.0f - dot(d, d)));
/* Create ray. */
@ -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. */

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

@ -752,19 +752,15 @@ float Camera::world_to_raster_size(float3 P)
/* No differentials, just use from directly ahead. */
camera_sample_panorama(&kernel_camera,
kernel_camera_motion.data(),
0.5f * full_width,
0.5f * full_height,
0.0f,
0.0f,
0.5f * make_float2(full_width, full_height),
zero_float2(),
&ray);
}
#else
camera_sample_panorama(&kernel_camera,
kernel_camera_motion.data(),
0.5f * full_width,
0.5f * full_height,
0.0f,
0.0f,
0.5f * make_float2(full_width, full_height),
zero_float2(),
&ray);
#endif

3
intern/cycles/scene/image_vdb.cpp
View File

@ -72,6 +72,9 @@ struct ToNanoOp {
catch (const std::exception &e) {
VLOG_WARNING << "Error converting OpenVDB to NanoVDB grid: " << e.what();
}
catch (...) {
VLOG_WARNING << "Error converting OpenVDB to NanoVDB grid: Unknown error";
}
return true;
}
else {

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;

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