Shadow linking: Initial work towards MIS support #107439
@ -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),
|
||||
|
@ -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. */
|
||||
|
@ -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:
|
||||
|
@ -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:
|
||||
|
@ -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;
|
||||
|
@ -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]);
|
||||
|
@ -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,
|
||||
|
@ -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];
|
||||
|
@ -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)
|
||||
|
@ -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
|
||||
|
@ -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
|
||||
|
@ -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)
|
||||
|
@ -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,
|
||||
|
@ -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,
|
||||
|
@ -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);
|
||||
|
@ -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);
|
||||
}
|
||||
|
@ -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
111
intern/cycles/kernel/integrator/intersect_dedicated_light.h
Normal 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
|
@ -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;
|
||||
}
|
||||
|
@ -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,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
124
intern/cycles/kernel/integrator/shade_dedicated_light.h
Normal 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
|
@ -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
|
||||
|
||||
|
@ -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
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
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(
|
||||
|
@ -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,
|
||||
Not sure what a good name is. Best I could think of so far is to rename
blocked_light
->light_dedicated
.