blender
/
blender
134
398
Fork 580
Code Issues 6.3k Pull Requests 757 Projects 19 Wiki Activity

Shadow linking: Initial work towards MIS support #107439

Merged
Sergey Sharybin merged 17 commits from Sergey/blender:cycles-light-linking-mis into cycles-light-linking 2023-05-05 21:44:14 +02:00
Conversation 8 Commits 17 Files Changed 33 +700 -147
33 changed files with 700 additions and 147 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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),
Brecht Van Lommel commented 2023-05-03 18:55:31 +02:00
Outdated
Review
Copy Link

Not sure what a good name is. Best I could think of so far is to rename blocked_light -> light_dedicated.

Not sure what a good name is. Best I could think of so far is to rename `blocked_light` -> `light_dedicated`.
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:

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

@ -38,6 +38,7 @@ bool kernel_has_intersection(DeviceKernel device_kernel)
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_INTERSECT_DEDICATED_LIGHT ||
device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE);
}
@ -46,6 +47,10 @@ struct ShaderCache {
ShaderCache(id<MTLDevice> _mtlDevice) : mtlDevice(_mtlDevice)
{
/* Initialize occupancy tuning LUT. */
// TODO: Look into tuning for DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT and
// DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT.
if (MetalInfo::get_device_vendor(mtlDevice) == METAL_GPU_APPLE) {
switch (MetalInfo::get_apple_gpu_architecture(mtlDevice)) {
default:

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

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

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

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

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

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

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

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

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

@ -502,7 +502,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);
@ -515,7 +516,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);
@ -1171,7 +1173,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

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

@ -119,6 +119,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:
@ -126,6 +127,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;
@ -409,6 +411,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);
@ -448,6 +459,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);

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;

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

@ -0,0 +1,111 @@
/* 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"
CCL_NAMESPACE_BEGIN
#ifdef __SHADOW_LINKING__
/* 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. */
ccl_device bool shadow_linking_pick_light_intersection(KernelGlobals kg,
IntegratorState state,
ccl_private const Ray *ccl_restrict ray,
ccl_private Intersection *ccl_restrict
isect)
{
uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
const int last_prim = INTEGRATOR_STATE(state, isect, prim);
const int last_object = INTEGRATOR_STATE(state, isect, object);
const int last_type = INTEGRATOR_STATE(state, isect, type);
/* The lights_intersect() has a "refining" behavior: it chooses intersection closer to the
* current intersection's distance. Hence initialize the fields which are accessed prior to
* recording an intersection. */
isect->t = FLT_MAX;
isect->prim = PRIM_NONE;
// TODO: Support mesh emitters.
// TODO: Support multiple light sources.
// TODO: Distant lights.
// TODO: Only if ray is not fully occluded.
// TODO: What of the actual shadow ray hits the same light through a semi-transparent surface?
if (!lights_intersect_shadow_linked(
kg, ray, isect, last_prim, last_object, last_type, path_flag)) {
return false;
}
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);
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

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

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;

22
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,8 @@ 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 (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;
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
continue;
}
#endif
/* Light linking. */

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

@ -0,0 +1,124 @@
/* 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"
CCL_NAMESPACE_BEGIN
#ifdef __SHADOW_LINKING__
/* 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)
{
kernel_assert(isect->type == PRIMITIVE_LAMP);
/* 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);
// TODO: Support mesh lights.
ray->self.light_object = OBJECT_NONE;
ray->self.light_prim = PRIM_NONE;
ray->self.light = isect->prim;
}
ccl_device void shadow_linking_shade(KernelGlobals kg, IntegratorState state)
{
/* 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);
LightSample ls ccl_optional_struct_init;
const bool use_light_sample = light_sample_from_intersection(kg, &isect, ray.P, ray.D, &ls);
if (!use_light_sample) {
/* No light to be sampled, so no direct light contribution either. */
return;
}
ShaderDataCausticsStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
const Spectrum light_eval = light_sample_shader_eval(kg, state, emission_sd, &ls, ray.time);
if (is_zero(light_eval)) {
return;
}
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
return;
}
/* MIS weighting. */
float mis_weight = 1.0f;
if (!(path_flag & PATH_RAY_MIS_SKIP)) {
mis_weight = light_sample_mis_weight_forward_lamp(kg, state, path_flag, &ls, ray.P);
}
const Spectrum bsdf_spectrum = light_eval * mis_weight;
shadow_linking_setup_ray_from_intersection(state, &ray, &isect);
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
kg, state, &ls, &ray, bsdf_spectrum, 0);
/* 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;
// TODO: Disable path guiding for this shadow ray?
}
#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);
#else
kernel_assert(!"integrator_intersect_dedicated_light is not supposed to be scheduled");
#endif
/* Restore self-intersection check primitives in the main state before returning to the
* intersect_closest() state. */
shadow_linking_restore_last_primitives(state);
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

209
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"
@ -135,6 +137,86 @@ 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 LightSample *ls,
ccl_private const Ray *ccl_restrict ray,
const Spectrum bsdf_spectrum,
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) = (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
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>
@ -243,12 +325,13 @@ 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);
Brecht Van Lommel commented 2023-05-05 20:38:33 +02:00
Review
Copy Link

It's a different type of offset. This one for float precision issues, the other one is for differences between smooth normals and actual geometry.

It's a different type of offset. This one for float precision issues, the other one is for differences between smooth normals and actual geometry.
Sergey Sharybin commented 2023-05-05 21:43:36 +02:00
Review
Copy Link

Ah, duuh. Now when you explained it, seems obvious. Thanks for explanation! :)

Ah, duuh. Now when you explained it, seems obvious. Thanks for explanation! :)
}
/* Copy volume stack and enter/exit volume. */
integrator_state_copy_volume_stack_to_shadow(kg, shadow_state, state);
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
kg, state, &ls, &ray, bsdf_eval_sum(&bsdf_eval), mnee_vertex_count);
if (is_transmission) {
#ifdef __VOLUME__
@ -256,18 +339,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);
integrator_state_write_shadow_ray_self(kg, shadow_state, &ray);
/* 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;
@ -289,55 +361,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. */
@ -588,9 +612,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);
@ -643,7 +667,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. */
@ -657,7 +681,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. */
@ -697,7 +721,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);
@ -710,7 +734,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,
@ -719,19 +756,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(

30
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"
@ -1213,27 +1214,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;
}
if (!kernel_data.integrator.use_light_mis) {
/* No need to cast extra shadow linking path if there are no lights with MIS in the scene. */
return false;
}
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

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

@ -135,3 +135,11 @@ 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)
/* 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)

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

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

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

@ -244,14 +244,16 @@ 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)
template<bool is_main_path>
ccl_device bool 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)
{
for (int lamp = 0; lamp < kernel_data.integrator.num_lights; lamp++) {
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
@ -269,9 +271,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
@ -283,10 +283,31 @@ 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) {
const bool is_indirect_ray = !(path_flag & PATH_RAY_CAMERA);
if (is_indirect_ray && kernel_data_fetch(lights, lamp).shadow_set_membership) {
continue;
}
}
else if (!kernel_data_fetch(lights, lamp).shadow_set_membership) {
continue;
}
}
#endif
#ifdef __LIGHT_LINKING__
/* Light linking. */
if (!light_link_light_match(kg, light_link_receiver_forward(kg, state), lamp)) {
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;
@ -325,6 +346,34 @@ ccl_device bool lights_intersect(KernelGlobals kg,
return isect->prim != PRIM_NONE;
}
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);
return lights_intersect_impl<true>(
kg, ray, isect, last_prim, last_object, last_type, path_flag, path_mnee, receiver_forward);
}
ccl_device bool 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)
{
return lights_intersect_impl<false>(
kg, ray, isect, last_prim, last_object, last_type, path_flag, PATH_MNEE_NONE, OBJECT_NONE);
}
/* Setup light sample from intersection. */
ccl_device bool light_sample_from_intersection(KernelGlobals kg,

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

@ -313,6 +313,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),
@ -1548,6 +1550,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,
@ -1555,6 +1558,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,
@ -1675,10 +1679,8 @@ enum KernelFeatureFlag : uint32_t {
/* OSL. */
KERNEL_FEATURE_OSL = (1U << 26U),
/* Light linking. */
/* Light and shadow linking. */
KERNEL_FEATURE_LIGHT_LINKING = (1U << 27U),
/* Shadow linking. */
KERNEL_FEATURE_SHADOW_LINKING = (1U << 28U),
};

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

@ -1093,12 +1093,6 @@ void LightManager::device_update_lights(Device *device, DeviceScene *dscene, Sce
klights[light_index].spot.spot_smooth = spot_smooth;
}
/* Disable MIS if the light participates in the shadow linking, as it is not supported. */
/* TODO(sergey): Support MIS with shadow linking. */
if (light->shadow_set_membership) {
shader_id &= ~SHADER_USE_MIS;
}
klights[light_index].shader_id = shader_id;
klights[light_index].max_bounces = max_bounces;

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

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

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

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