ilusha/blender
1
0
Fork 0
mirror of https://github.com/blender/blender synced 2026-09-29 04:37:17 +03:00
Code Issues Projects Releases Packages Wiki Activity Actions

Cycles: Perform direct light shader eval in own kernel

Browse source 
This improves performance by 5-10% for various benchmark scenes and GPU
devices, while on others it's roughly the same. There is a performance
regression with Intel Arc A750 on Linux related to shadow queueing
overhead, that is planned to be fixed separately.

Another goal of this change is to sidestep GPU compiler bugs that seems
more likely to happen with bigger kernels, and to make it easier for the
texture cache to cancel and resume on cache miss.

A new shade_light_nee kernel was added, and shade_light was renamed to
shade_light_forward (following naming for MIS functions). The shade_light_nee
kernel is only used when the light does not have constant emission.

The shade_dedicate_light kernel no longer does any shading. A future
optimization may be to fold this into the intersect_dedicated_light kernel.

LightSample.uv was removed as shading no longer happens immediately. A new
LightPdf was added for the cases where only the pdf is needed, avoiding the
overhead of constructing a full LightSample. There may be more room to
shrink LightSample in future refactors.

The integrate state memory usage is increased by 1 float when not using the
light tree, for the light threshold. All other informating for shading is
reconstructed the shadow ray, including position, normal and uv.

Pull Request: https://projects.blender.org/blender/blender/pulls/152649
This commit is contained in:
Brecht Van Lommel 2026-01-07 18:34:33 +01:00
parent fe2b3681aa
commit 4b34743b4e
32 changed files with 734 additions and 479 deletions
Download patch file Download diff file Expand all files Collapse all files

9
intern/cycles/device/kernel.cpp
View file

@ -13,7 +13,8 @@ CCL_NAMESPACE_BEGIN
bool device_kernel_has_shading(DeviceKernel kernel)
{
return (kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE ||
@ -58,8 +59,10 @@ const char *device_kernel_as_string(DeviceKernel kernel)
return "integrator_intersect_dedicated_light";
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
return "integrator_shade_background";
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
return "integrator_shade_light";
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE:
return "integrator_shade_light_nee";
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD:
return "integrator_shade_light_forward";
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
return "integrator_shade_shadow";
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:

2
intern/cycles/device/metal/kernel.mm
View file

@ -46,7 +46,7 @@ struct ShaderCache {
/* Initialize occupancy tuning LUT. */
// TODO: Look into tuning for DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT and
// DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT.
// DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_*.
switch (MetalInfo::get_apple_gpu_architecture(mtlDevice)) {
default:

3
intern/cycles/device/oneapi/device_impl.cpp
View file

@ -1317,7 +1317,8 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:

15
intern/cycles/device/optix/device_impl.cpp
View file

@ -558,10 +558,14 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
group_descs[PG_RGEN_SHADE_BACKGROUND].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_BACKGROUND].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_shade_background";
group_descs[PG_RGEN_SHADE_LIGHT].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_SHADE_LIGHT].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_LIGHT].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_shade_light";
group_descs[PG_RGEN_SHADE_LIGHT_NEE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_SHADE_LIGHT_NEE].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_LIGHT_NEE].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_shade_light_nee";
group_descs[PG_RGEN_SHADE_LIGHT_FORWARD].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_SHADE_LIGHT_FORWARD].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_LIGHT_FORWARD].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_shade_light_forward";
group_descs[PG_RGEN_SHADE_SURFACE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_SHADE_SURFACE].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_SURFACE].raygen.entryFunctionName =
@ -1052,7 +1056,8 @@ bool OptiXDevice::load_osl_kernels()
vector<OptixProgramGroup> pipeline_groups;
pipeline_groups.reserve(NUM_PROGRAM_GROUPS);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_BACKGROUND]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_LIGHT]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_LIGHT_NEE]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_LIGHT_FORWARD]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_RAYTRACE]);
pipeline_groups.push_back(groups[PG_CALL_SVM_AO]);

3
intern/cycles/device/optix/device_impl.h
View file

@ -27,7 +27,8 @@ enum {
PG_RGEN_INTERSECT_VOLUME_STACK,
PG_RGEN_INTERSECT_DEDICATED_LIGHT,
PG_RGEN_SHADE_BACKGROUND,
PG_RGEN_SHADE_LIGHT,
PG_RGEN_SHADE_LIGHT_NEE,
PG_RGEN_SHADE_LIGHT_FORWARD,
PG_RGEN_SHADE_SURFACE,
PG_RGEN_SHADE_SURFACE_RAYTRACE,
PG_RGEN_SHADE_SURFACE_MNEE,

8
intern/cycles/device/optix/queue.cpp
View file

@ -104,9 +104,13 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel,
pipeline = optix_device->pipelines[PIP_SHADE];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_BACKGROUND * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE:
pipeline = optix_device->pipelines[PIP_SHADE];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_LIGHT * sizeof(SbtRecord);
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_LIGHT_NEE * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD:
pipeline = optix_device->pipelines[PIP_SHADE];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_LIGHT_FORWARD * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
pipeline = optix_device->pipelines[PIP_SHADE];

11
intern/cycles/integrator/path_trace_work_gpu.cpp
View file

@ -476,6 +476,10 @@ bool PathTraceWorkGPU::enqueue_path_iteration()
const int available_shadow_paths = max_num_paths_ -
integrator_next_shadow_path_index_.data()[0];
if (available_shadow_paths < queue_counter->num_queued[kernel]) {
if (queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE]) {
enqueue_path_iteration(DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE);
return true;
}
if (queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW]) {
enqueue_path_iteration(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW);
return true;
@ -558,7 +562,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
@ -690,6 +695,7 @@ void PathTraceWorkGPU::compact_shadow_paths()
{
IntegratorQueueCounter *queue_counter = integrator_queue_counter_.data();
const int num_active_paths =
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE] +
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] +
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW];
@ -1289,7 +1295,8 @@ bool PathTraceWorkGPU::kernel_creates_ao_paths(DeviceKernel kernel)
bool PathTraceWorkGPU::kernel_is_shadow_path(DeviceKernel kernel)
{
return (kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW);
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE);
}
int PathTraceWorkGPU::kernel_max_active_main_path_index(DeviceKernel kernel)

19
intern/cycles/kernel/device/gpu/kernel.h
View file

@ -235,7 +235,7 @@ 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_light,
ccl_gpu_kernel_signature(integrator_shade_light_nee,
const ccl_global int *path_index_array,
ccl_global float *render_buffer,
const int work_size)
@ -244,7 +244,22 @@ ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
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_light(nullptr, state, render_buffer));
ccl_gpu_kernel_call(integrator_shade_light_nee(nullptr, state, render_buffer));
}
}
ccl_gpu_kernel_postfix
ccl_gpu_kernel(GPU_KERNEL_BLOCK_NUM_THREADS, GPU_KERNEL_MAX_REGISTERS)
ccl_gpu_kernel_signature(integrator_shade_light_forward,
const ccl_global 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_light_forward(nullptr, state, render_buffer));
}
}
ccl_gpu_kernel_postfix

13
intern/cycles/kernel/device/oneapi/kernel.cpp
View file

@ -432,9 +432,18 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_background);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT: {
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE: {
oneapi_call(
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_light);
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_light_nee);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD: {
oneapi_call(kg,
cgh,
global_size,
local_size,
args,
oneapi_kernel_integrator_shade_light_forward);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW: {

13
intern/cycles/kernel/device/optix/kernel_osl.cu
View file

@ -26,13 +26,22 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_shade_background()
integrator_shade_background(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_light()
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_light_nee()
{
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_light(nullptr, path_index, kernel_params.render_buffer);
integrator_shade_light_nee(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_light_forward()
{
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_light_forward(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface()

4
intern/cycles/kernel/integrator/guiding.h
View file

@ -447,6 +447,10 @@ ccl_device_forceinline void guiding_record_direct_light(KernelGlobals kg,
if (!kernel_data.integrator.train_guiding) {
return;
}
const uint32_t path_flag = INTEGRATOR_STATE(state, shadow_path, flag);
if (path_flag & PATH_RAY_SHADOW_FOR_AO) {
return;
}
if (state->shadow_path.path_segment) {
const Spectrum Lo = safe_divide_color(INTEGRATOR_STATE(state, shadow_path, throughput),
INTEGRATOR_STATE(state, shadow_path, unlit_throughput));

4
intern/cycles/kernel/integrator/intersect_closest.h
View file

@ -251,7 +251,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel(
if (hit) {
/* Hit a surface, continue with light or surface kernel. */
if (isect->type & PRIMITIVE_LAMP) {
integrator_path_next(state, current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT);
integrator_path_next(state, current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD);
}
else {
/* Hit a surface, continue with surface kernel unless terminated. */
@ -311,7 +311,7 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_volume(
if (isect->prim != PRIM_NONE) {
/* Hit a surface, continue with light or surface kernel. */
if (isect->type & PRIMITIVE_LAMP) {
integrator_path_next(state, current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT);
integrator_path_next(state, current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD);
return;
}

7
intern/cycles/kernel/integrator/megakernel.h
View file

@ -36,6 +36,9 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
integrator_shade_shadow(kg, &state->shadow, render_buffer);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE:
integrator_shade_light_nee(kg, &state->shadow, render_buffer);
break;
default:
kernel_assert(0);
break;
@ -85,8 +88,8 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
integrator_shade_surface_mnee(kg, state, render_buffer);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT:
integrator_shade_light(kg, state, render_buffer);
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD:
integrator_shade_light_forward(kg, state, render_buffer);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT:
integrator_shade_dedicated_light(kg, state, render_buffer);

12
intern/cycles/kernel/integrator/mnee.h
View file

@ -826,24 +826,14 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg,
/* Set diffuse bounce info. */
INTEGRATOR_STATE_WRITE(state, path, diffuse_bounce) = diffuse_bounce + 1;
/* Evaluate light sample
* in case the light has a node-based shader:
* 1. sd_mnee will be used to store light data, which is why we need to do
* this evaluation here. sd_mnee needs to contain the solution's last
* interface data at the end of the call for the shadow ray setup to work.
* 2. ls needs to contain the last interface data for the light shader to
* evaluate properly */
/* Set bounce info in case a light path node is used in the light shader graph. */
INTEGRATOR_STATE_WRITE(state, path, transmission_bounce) = transmission_bounce + vertex_count -
1;
INTEGRATOR_STATE_WRITE(state, path, bounce) = bounce + vertex_count;
const Spectrum light_eval = light_sample_shader_eval(kg, state, sd_mnee, ls, sd->time);
bsdf_eval_mul(throughput, light_eval / ls->pdf);
bsdf_eval_mul(throughput, ls->eval_fac / ls->pdf);
/* Generalized geometry term. */
float dh_dx;
float dx1_dxlight;
if (!mnee_compute_transfer_matrix(

90
intern/cycles/kernel/integrator/shade_background.h
View file

@ -14,8 +14,11 @@
#include "kernel/light/light.h"
#include "kernel/light/sample.h"
#include "kernel/geom/object.h"
#include "kernel/geom/shader_data.h"
#include "kernel/types.h"
CCL_NAMESPACE_BEGIN
ccl_device Spectrum integrator_eval_background_shader(KernelGlobals kg,
@ -127,59 +130,66 @@ ccl_device_inline void integrate_distant_lights(KernelGlobals kg,
{
const float3 ray_D = INTEGRATOR_STATE(state, ray, D);
const float ray_time = INTEGRATOR_STATE(state, ray, time);
LightSample ls ccl_optional_struct_init;
for (int lamp = 0; lamp < kernel_data.integrator.num_lights; lamp++) {
if (distant_light_sample_from_intersection(kg, ray_D, lamp, &ls)) {
/* Use visibility flag to skip lights. */
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
if (klight->type != LIGHT_DISTANT || !(klight->shader_id & SHADER_USE_MIS)) {
continue;
}
LightEval light_eval = distant_light_eval_from_intersection(klight, ray_D);
if (light_eval.eval_fac == 0.0f) {
continue;
}
/* Use visibility flag to skip lights. */
#ifdef __PASSES__
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
continue;
}
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
if (!is_light_shader_visible_to_path(klight->shader_id, path_flag)) {
continue;
}
#endif
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
#ifdef __LIGHT_LINKING__
if (!light_link_light_match(kg, light_link_receiver_forward(kg, state), klight->object_id) &&
!(path_flag & PATH_RAY_CAMERA))
{
continue;
}
if (!light_link_light_match(kg, light_link_receiver_forward(kg, state), klight->object_id) &&
!(path_flag & PATH_RAY_CAMERA))
{
continue;
}
#endif
#ifdef __SHADOW_LINKING__
if (kernel_data_fetch(objects, klight->object_id).shadow_set_membership !=
LIGHT_LINK_MASK_ALL)
{
continue;
}
if (kernel_data_fetch(objects, klight->object_id).shadow_set_membership != LIGHT_LINK_MASK_ALL)
{
continue;
}
#endif
#ifdef __MNEE__
if (INTEGRATOR_STATE(state, path, mnee) & PATH_MNEE_CULL_LIGHT_CONNECTION) {
/* This path should have been resolved with mnee, it will
* generate a firefly for small lights since it is improbable. */
if (klight->use_caustics) {
continue;
}
}
#endif /* __MNEE__ */
/* Evaluate light shader. */
/* TODO: does aliasing like this break automatic SoA in CUDA? */
ShaderDataTinyStorage 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)) {
if (INTEGRATOR_STATE(state, path, mnee) & PATH_MNEE_CULL_LIGHT_CONNECTION) {
/* This path should have been resolved with mnee, it will
* generate a firefly for small lights since it is improbable. */
if (klight->use_caustics) {
continue;
}
/* MIS weighting. */
const float mis_weight = light_sample_mis_weight_forward_distant(kg, state, path_flag, &ls);
/* Write to render buffer. */
guiding_record_background(kg, state, light_eval, mis_weight);
film_write_surface_emission(kg, state, light_eval, mis_weight, render_buffer, ls.group);
}
#endif /* __MNEE__ */
/* Evaluate light shader. */
const Spectrum shader_eval = light_sample_shader_eval_forward(
kg, state, lamp, zero_float3(), ray_D, FLT_MAX, ray_time);
const float3 eval = shader_eval * light_eval.eval_fac;
if (is_zero(eval)) {
continue;
}
/* MIS weighting. */
const float mis_weight = light_sample_mis_weight_forward_distant(
kg, state, path_flag, lamp, light_eval.pdf);
/* Write to render buffer. */
guiding_record_background(kg, state, eval, mis_weight);
film_write_surface_emission(
kg, state, eval, mis_weight, render_buffer, object_lightgroup(kg, klight->object_id));
}
}

116
intern/cycles/kernel/integrator/shade_dedicated_light.h
View file

@ -14,37 +14,34 @@ CCL_NAMESPACE_BEGIN
#ifdef __SHADOW_LINKING__
ccl_device_inline bool shadow_linking_light_sample_from_intersection(
KernelGlobals kg,
const ccl_private Intersection &ccl_restrict isect,
const ccl_private Ray &ccl_restrict ray,
const float3 N,
const uint32_t path_flag,
ccl_private LightSample *ccl_restrict ls)
ccl_device_inline LightEval
shadow_linking_light_eval_from_intersection(KernelGlobals kg,
const ccl_private Intersection &ccl_restrict isect,
const ccl_private Ray &ccl_restrict ray,
const float3 N,
const uint32_t path_flag)
{
const int lamp = isect.prim;
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, isect.prim);
const LightType type = LightType(klight->type);
if (type == LIGHT_DISTANT) {
return distant_light_sample_from_intersection(kg, ray.D, lamp, ls);
}
return light_sample_from_intersection(kg, &isect, ray.P, ray.D, N, path_flag, ls);
return (type == LIGHT_DISTANT) ?
distant_light_eval_from_intersection(klight, ray.D) :
light_eval_from_intersection(kg, &isect, ray.P, ray.D, N, path_flag);
}
ccl_device_inline float shadow_linking_light_sample_mis_weight(KernelGlobals kg,
IntegratorState state,
const uint32_t path_flag,
const ccl_private LightSample *ls,
const int light_id,
const float light_sample_pdf,
const float3 P)
{
if (ls->type == LIGHT_DISTANT) {
return light_sample_mis_weight_forward_distant(kg, state, path_flag, ls);
if (kernel_data_fetch(lights, light_id).type == LIGHT_DISTANT) {
return light_sample_mis_weight_forward_distant(
kg, state, path_flag, light_id, light_sample_pdf);
}
return light_sample_mis_weight_forward_lamp(kg, state, path_flag, ls, P);
return light_sample_mis_weight_forward_lamp(kg, state, path_flag, light_id, light_sample_pdf, P);
}
/* Setup ray for the shadow path.
@ -75,49 +72,48 @@ ccl_device bool shadow_linking_shade_light(KernelGlobals kg,
IntegratorState state,
ccl_private Ray &ccl_restrict ray,
ccl_private Intersection &ccl_restrict isect,
ccl_private ShaderData *emission_sd,
ccl_private Spectrum &ccl_restrict bsdf_spectrum,
ccl_private float &ccl_restrict light_weight,
ccl_private float &mis_weight,
ccl_private int &ccl_restrict light_group)
ccl_private int &ccl_restrict light_group,
ccl_private int &ccl_restrict shader_id)
{
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
const float3 N = INTEGRATOR_STATE(state, path, mis_origin_n);
LightSample ls ccl_optional_struct_init;
const bool use_light_sample = shadow_linking_light_sample_from_intersection(
kg, isect, ray, N, path_flag, &ls);
if (!use_light_sample) {
const LightEval light_eval = shadow_linking_light_eval_from_intersection(
kg, isect, ray, N, path_flag);
if (light_eval.eval_fac == 0.0f) {
/* No light to be sampled, so no direct light contribution either. */
return false;
}
const Spectrum light_eval = light_sample_shader_eval(kg, state, emission_sd, &ls, ray.time);
if (is_zero(light_eval)) {
return false;
}
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, isect.prim);
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
if (!is_light_shader_visible_to_path(klight->shader_id, path_flag)) {
return false;
}
/* MIS weighting. */
mis_weight = shadow_linking_light_sample_mis_weight(kg, state, path_flag, &ls, ray.P);
mis_weight = shadow_linking_light_sample_mis_weight(
kg, state, path_flag, isect.prim, light_eval.pdf, ray.P);
bsdf_spectrum = light_eval * mis_weight *
INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
light_group = ls.group;
light_weight = light_eval.eval_fac * mis_weight *
INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
light_group = object_lightgroup(kg, klight->object_id);
shader_id = klight->shader_id;
return true;
}
ccl_device bool shadow_linking_shade_surface_emission(KernelGlobals kg,
IntegratorState state,
ccl_private ShaderData *emission_sd,
ccl_global float *ccl_restrict render_buffer,
ccl_private Spectrum &ccl_restrict
bsdf_spectrum,
ccl_private float &ccl_restrict light_weight,
ccl_private float &mis_weight,
ccl_private int &ccl_restrict light_group)
ccl_private int &ccl_restrict light_group,
ccl_private int &ccl_restrict shader_id)
{
ShaderDataTinyStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
integrate_surface_shader_setup(kg, state, emission_sd);
@ -128,26 +124,16 @@ ccl_device bool shadow_linking_shade_surface_emission(KernelGlobals kg,
}
# endif
surface_shader_eval<KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT>(
kg, state, emission_sd, render_buffer, path_flag | PATH_RAY_EMISSION);
if ((emission_sd->flag & SD_EMISSION) == 0) {
return false;
}
const Spectrum L = surface_shader_emission(emission_sd);
mis_weight = light_sample_mis_weight_forward_surface(kg, state, path_flag, emission_sd);
bsdf_spectrum = L * mis_weight * INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
light_weight = mis_weight * INTEGRATOR_STATE(state, shadow_link, dedicated_light_weight);
light_group = object_lightgroup(kg, emission_sd->object);
shader_id = emission_sd->shader;
return true;
}
ccl_device void shadow_linking_shade(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
ccl_device void shadow_linking_shade(KernelGlobals kg, IntegratorState state)
{
/* Read intersection from integrator state into local memory. */
Intersection isect ccl_optional_struct_init;
@ -157,29 +143,33 @@ ccl_device void shadow_linking_shade(KernelGlobals kg,
Ray ray ccl_optional_struct_init;
integrator_state_read_ray(state, &ray);
ShaderDataTinyStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
Spectrum bsdf_spectrum;
float light_weight = 0.0f;
float mis_weight = 1.0f;
int light_group = LIGHTGROUP_NONE;
int shader_id = SHADER_NONE;
if (isect.type == PRIMITIVE_LAMP) {
if (!shadow_linking_shade_light(
kg, state, ray, isect, emission_sd, bsdf_spectrum, mis_weight, light_group))
kg, state, ray, isect, light_weight, mis_weight, light_group, shader_id))
{
return;
}
}
else {
if (!shadow_linking_shade_surface_emission(
kg, state, emission_sd, render_buffer, bsdf_spectrum, mis_weight, light_group))
kg, state, light_weight, mis_weight, light_group, shader_id))
{
return;
}
}
if (is_zero(bsdf_spectrum)) {
/* Evaluate constant part of light shader, rest will optionally be done in another kernel. */
Spectrum light_eval;
const bool is_constant_light_shader = light_sample_shader_eval_nee_constant(
kg, shader_id, isect.prim, isect.type == PRIMITIVE_LAMP, light_eval);
light_eval *= light_weight;
if (is_zero(light_eval)) {
return;
}
@ -187,7 +177,7 @@ ccl_device void shadow_linking_shade(KernelGlobals kg,
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
kg, state, &ray, bsdf_spectrum, light_group, 0);
kg, state, &ray, light_eval, light_group, 0, is_constant_light_shader);
/* The light is accumulated from the shade_surface kernel, which will make the clamping decision
* based on the actual value of the bounce. For the dedicated shadow ray we want to follow the
@ -223,12 +213,12 @@ ccl_device void shadow_linking_shade(KernelGlobals kg,
ccl_device void integrator_shade_dedicated_light(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
ccl_global float *ccl_restrict /*render_buffer*/)
{
PROFILING_INIT(kg, PROFILING_SHADE_DEDICATED_LIGHT);
#ifdef __SHADOW_LINKING__
shadow_linking_shade(kg, state, render_buffer);
shadow_linking_shade(kg, state);
/* Restore self-intersection check primitives in the main state before returning to the
* intersect_closest() state. */

180
intern/cycles/kernel/integrator/shade_light.h
View file

@ -6,14 +6,19 @@
#include "kernel/film/light_passes.h"
#include "kernel/integrator/path_state.h"
#include "kernel/light/light.h"
#include "kernel/light/sample.h"
#include "kernel/geom/object.h"
#include "kernel/types.h"
CCL_NAMESPACE_BEGIN
ccl_device_inline void integrate_light(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
ccl_device_inline void integrate_light_forward(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
{
/* Setup light sample. */
Intersection isect ccl_optional_struct_init;
@ -30,45 +35,49 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
/* Advance ray to new start distance. */
INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(isect.t);
LightSample ls ccl_optional_struct_init;
const bool use_light_sample = light_sample_from_intersection(
kg, &isect, ray_P, ray_D, N, path_flag, &ls);
if (!use_light_sample) {
const LightEval light_eval = light_eval_from_intersection(
kg, &isect, ray_P, ray_D, N, path_flag);
if (light_eval.eval_fac == 0.0f) {
return;
}
/* Use visibility flag to skip lights. */
#ifdef __PASSES__
if (!is_light_shader_visible_to_path(ls.shader, path_flag)) {
return;
{
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, isect.prim);
if (!is_light_shader_visible_to_path(klight->shader_id, path_flag)) {
return;
}
}
#endif
/* Evaluate light shader. */
/* TODO: does aliasing like this break automatic SoA in CUDA? */
ShaderDataTinyStorage 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)) {
const Spectrum shader_eval = light_sample_shader_eval_forward(
kg, state, isect.prim, ray_P, ray_D, isect.t, ray_time);
const float3 eval = shader_eval * light_eval.eval_fac;
if (is_zero(eval)) {
return;
}
/* MIS weighting. */
const float mis_weight = light_sample_mis_weight_forward_lamp(kg, state, path_flag, &ls, ray_P);
const float mis_weight = light_sample_mis_weight_forward_lamp(
kg, state, path_flag, isect.prim, light_eval.pdf, ray_P);
/* Write to render buffer. */
guiding_record_surface_emission(kg, state, light_eval, mis_weight);
film_write_surface_emission(kg, state, light_eval, mis_weight, render_buffer, ls.group);
guiding_record_surface_emission(kg, state, eval, mis_weight);
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, isect.prim);
film_write_surface_emission(
kg, state, eval, mis_weight, render_buffer, object_lightgroup(kg, klight->object_id));
}
ccl_device void integrator_shade_light(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
/* Evaluate light shader at intersection in forward path tracing. */
ccl_device void integrator_shade_light_forward(KernelGlobals kg,
IntegratorState state,
ccl_global float *ccl_restrict render_buffer)
{
PROFILING_INIT(kg, PROFILING_SHADE_LIGHT_SETUP);
integrate_light(kg, state, render_buffer);
integrate_light_forward(kg, state, render_buffer);
/* TODO: we could get stuck in an infinite loop if there are precision issues
* and the same light is hit again.
@ -80,16 +89,137 @@ ccl_device void integrator_shade_light(KernelGlobals kg,
INTEGRATOR_STATE_WRITE(state, path, transparent_bounce) = transparent_bounce;
if (transparent_bounce >= kernel_data.integrator.transparent_max_bounce) {
integrator_path_terminate(kg, state, render_buffer, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT);
integrator_path_terminate(
kg, state, render_buffer, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD);
return;
}
integrator_path_next(
state, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
integrator_path_next(state,
DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
/* TODO: in some cases we could continue directly to SHADE_BACKGROUND, but
* probably that optimization is probably not practical if we add lights to
* scene geometry. */
}
ccl_device bool integrate_light_nee(KernelGlobals kg, IntegratorShadowState state)
{
/* Read intersection and ray. */
Ray ray ccl_optional_struct_init;
integrator_state_read_shadow_ray(state, &ray);
integrator_state_read_shadow_ray_self(state, &ray);
Intersection isect = {};
isect.object = ray.self.light_object;
isect.prim = ray.self.light_prim;
isect.type = kernel_data_fetch(objects, isect.object).primitive_type;
isect.t = ray.tmax;
kernel_assert(isect.object != OBJECT_NONE);
kernel_assert(isect.prim != PRIM_NONE);
float3 eval = zero_spectrum();
bool is_background = false;
/* Setup shader data */
ShaderDataCausticsStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
PROFILING_INIT_FOR_SHADER(kg, PROFILING_SHADE_LIGHT_SETUP);
if (isect.type == PRIMITIVE_LAMP) {
/* Lights. */
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, isect.prim);
const LightType light_type = LightType(klight->type);
if (light_type == LIGHT_BACKGROUND) {
/* Background light. */
shader_setup_from_background(kg, emission_sd, ray.P, ray.D, ray.time);
is_background = true;
}
else {
/* Other light types.
* Compute Ng and UV on demand so we don't have to store it in integrator state. */
const float3 P = (ray.tmax == FLT_MAX) ? -ray.D : ray.P + ray.tmax * ray.D;
float3 Ng = zero_float3();
float2 uv = zero_float2();
light_normal_uv_from_position(kg, klight, P, ray.D, Ng, uv);
shader_setup_from_sample(kg,
emission_sd,
P,
Ng,
-ray.D,
klight->shader_id,
isect.object,
isect.prim,
uv.x,
uv.y,
ray.tmax,
ray.time,
false,
true);
}
}
else {
/* Triangles.
* Compute UV on demand so we don't have to store it in integrator state. */
const float2 uv = triangle_light_uv(kg, isect.object, isect.prim, ray.time, ray.P, ray.D);
isect.u = uv.x;
isect.v = uv.y;
shader_setup_from_ray(kg, emission_sd, &ray, &isect);
}
/* Evaluate shader. */
PROFILING_SHADER(emission_sd->object, emission_sd->shader);
PROFILING_EVENT(PROFILING_SHADE_LIGHT_EVAL);
/* No proper path flag, we're evaluating this for all closures. that's
* weak but we'd have to do multiple evaluations otherwise. */
surface_shader_eval<KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT>(
kg, state, emission_sd, nullptr, PATH_RAY_EMISSION);
/* Evaluate emission closures. */
eval = (is_background) ? surface_shader_background(emission_sd) :
surface_shader_emission(emission_sd);
/* Probabilistic light termination.
* Light threshold is only used without light tree. */
if (!(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_TREE)) {
RNGState rng_state;
shadow_path_state_rng_load(state, &rng_state);
const float rand_terminate = path_state_rng_light_termination(kg, &rng_state);
const float bsdf_eval_average = INTEGRATOR_STATE(state, shadow_path, bsdf_eval_average);
if (light_sample_terminate(kg, eval, bsdf_eval_average, rand_terminate)) {
return false;
}
}
else if (is_zero(eval)) {
return false;
}
/* Update throughput. */
INTEGRATOR_STATE(state, shadow_path, throughput) *= eval;
return true;
}
/* Evaluate light shader for next event estimation, after shade_surface and shade_volume and before
* shadow ray intersection. Only when the light has non-constant emisison. */
ccl_device void integrator_shade_light_nee(KernelGlobals kg,
IntegratorShadowState state,
ccl_global float *ccl_restrict /*render_buffer*/)
{
PROFILING_INIT(kg, PROFILING_SHADE_LIGHT_SETUP);
if (!integrate_light_nee(kg, state)) {
integrator_shadow_path_terminate(state, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE);
return;
}
integrator_shadow_path_next(
state, DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW);
}
CCL_NAMESPACE_END

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

@ -10,6 +10,7 @@
#include "kernel/integrator/volume_stack.h"
#include "kernel/geom/shader_data.h"
#include "kernel/light/light.h"
CCL_NAMESPACE_BEGIN

73
intern/cycles/kernel/integrator/shade_surface.h
View file

@ -208,12 +208,17 @@ integrate_direct_light_shadow_init_common(KernelGlobals kg,
const ccl_private Ray *ccl_restrict ray,
const Spectrum bsdf_spectrum,
const int light_group,
const int mnee_vertex_count)
const int mnee_vertex_count,
const bool constant_light_shader)
{
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
kg,
state,
(constant_light_shader) ? DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW :
DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE,
false);
#ifdef __VOLUME__
/* Copy volume stack and enter/exit volume. */
@ -228,6 +233,10 @@ integrate_direct_light_shadow_init_common(KernelGlobals kg,
const Spectrum unlit_throughput = INTEGRATOR_STATE(state, path, throughput);
const Spectrum throughput = unlit_throughput * bsdf_spectrum;
if (!(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_TREE)) {
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bsdf_eval_average) = average(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(
@ -341,15 +350,6 @@ ccl_device
}
}
/* Evaluate light shader.
*
* TODO: can we reuse sd memory? In theory we can move this after
* integrate_surface_bounce, evaluate the BSDF, and only then evaluate
* the light shader. This could also move to its own kernel, for
* non-constant light sources. */
ShaderDataCausticsStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
Ray ray ccl_optional_struct_init;
BsdfEval bsdf_eval ccl_optional_struct_init;
@ -368,34 +368,51 @@ ccl_device
/* Are we on a caustic receiver? */
if (!is_transmission && (sd->object_flag & SD_OBJECT_CAUSTICS_RECEIVER)) {
ShaderDataCausticsStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
mnee_vertex_count = kernel_path_mnee_sample(
kg, state, sd, emission_sd, rng_state, &ls, &bsdf_eval);
if (mnee_vertex_count > 0) {
/* Create shadow ray after successful manifold walk:
* emission_sd contains the last interface intersection and
* the light sample ls has been updated */
light_sample_to_surface_shadow_ray(kg, emission_sd, &ls, &ray);
}
}
}
}
}
#endif
/* Evaluate constant part of light shader, rest will optionally be done in another kernel. */
Spectrum light_shader_eval ccl_optional_struct_init;
const bool is_constant_light_shader = light_sample_shader_eval_nee_constant(
kg, ls.shader, ls.prim, ls.type != LIGHT_TRIANGLE, light_shader_eval);
#ifdef __MNEE__
if (mnee_vertex_count > 0) {
/* Create shadow ray after successful manifold walk:
* emission_sd contains the last interface intersection and
* the light sample ls has been updated */
light_sample_to_surface_shadow_ray(kg, emission_sd, &ls, &ray);
bsdf_eval_mul(&bsdf_eval, light_shader_eval);
}
else
#endif /* __MNEE__ */
{
const Spectrum light_eval = light_sample_shader_eval(kg, state, emission_sd, &ls, sd->time);
if (is_zero(light_eval)) {
return;
}
/* Evaluate BSDF. */
const float bsdf_pdf = surface_shader_bsdf_eval(kg, state, sd, ls.D, &bsdf_eval, ls.shader);
const float mis_weight = light_sample_mis_weight_nee(kg, ls.pdf, bsdf_pdf);
bsdf_eval_mul(&bsdf_eval, light_eval / ls.pdf * mis_weight);
bsdf_eval_mul(&bsdf_eval, light_shader_eval * ls.eval_fac / ls.pdf * mis_weight);
/* Path termination. */
const float terminate = path_state_rng_light_termination(kg, rng_state);
if (light_sample_terminate(kg, &bsdf_eval, terminate)) {
/* Path termination for constant light shader. */
if (is_constant_light_shader && !(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_TREE)) {
const float terminate = path_state_rng_light_termination(kg, rng_state);
if (light_sample_terminate(kg, &bsdf_eval, terminate)) {
return;
}
}
/* For non-constant light shader, probabilistic termination happens in
* SHADE_LIGHT_NEE when the full contribution is known. */
else if (bsdf_eval_is_zero(&bsdf_eval)) {
return;
}
@ -409,7 +426,13 @@ ccl_device
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrate_direct_light_shadow_init_common(
kg, state, &ray, bsdf_eval_sum(&bsdf_eval), ls.group, mnee_vertex_count);
kg,
state,
&ray,
bsdf_eval_sum(&bsdf_eval),
ls.group,
mnee_vertex_count,
is_constant_light_shader);
if (is_transmission) {
#ifdef __VOLUME__

45
intern/cycles/kernel/integrator/shade_volume.h
View file

@ -13,6 +13,7 @@
#include "kernel/integrator/intersect_closest.h"
#include "kernel/integrator/path_state.h"
#include "kernel/integrator/shadow_linking.h"
#include "kernel/integrator/state.h"
#include "kernel/integrator/volume_shader.h"
#include "kernel/integrator/volume_stack.h"
@ -2421,29 +2422,28 @@ ccl_device_forceinline void integrate_volume_direct_light(
return;
}
/* Evaluate light shader.
*
* TODO: can we reuse sd memory? In theory we can move this after
* integrate_surface_bounce, evaluate the BSDF, and only then evaluate
* the light shader. This could also move to its own kernel, for
* non-constant light sources. */
ShaderDataTinyStorage 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, sd->time);
if (is_zero(light_eval)) {
return;
}
/* Evaluate constant part of light shader, rest will optionally be done in another kernel. */
Spectrum light_shader_eval ccl_optional_struct_init;
const bool is_constant_light_shader = light_sample_shader_eval_nee_constant(
kg, ls.shader, ls.prim, ls.type != LIGHT_TRIANGLE, light_shader_eval);
/* Evaluate BSDF. */
BsdfEval phase_eval ccl_optional_struct_init;
const float phase_pdf = volume_shader_phase_eval(
kg, state, sd, phases, ls.D, &phase_eval, ls.shader);
const float mis_weight = light_sample_mis_weight_nee(kg, ls.pdf, phase_pdf);
bsdf_eval_mul(&phase_eval, light_eval / ls.pdf * mis_weight);
bsdf_eval_mul(&phase_eval, light_shader_eval * ls.eval_fac / ls.pdf * mis_weight);
/* Path termination. */
const float terminate = path_state_rng_light_termination(kg, rng_state);
if (light_sample_terminate(kg, &phase_eval, terminate)) {
/* Path termination for constant light shader. */
if (is_constant_light_shader && !(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_TREE)) {
const float terminate = path_state_rng_light_termination(kg, rng_state);
if (light_sample_terminate(kg, &phase_eval, terminate)) {
return;
}
}
/* For non-constant light shader, probablistic termination happens in
* SHADE_LIGHT_NEE when the full contribution is known. */
else if (bsdf_eval_is_zero(&phase_eval)) {
return;
}
@ -2453,7 +2453,11 @@ ccl_device_forceinline void integrate_volume_direct_light(
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg, state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW, false);
kg,
state,
(is_constant_light_shader) ? DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW :
DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE,
false);
/* Write shadow ray and associated state to global memory. */
integrator_state_write_shadow_ray(shadow_state, &ray);
@ -2463,7 +2467,12 @@ ccl_device_forceinline void integrate_volume_direct_light(
const uint16_t bounce = INTEGRATOR_STATE(state, path, bounce);
const uint16_t transparent_bounce = INTEGRATOR_STATE(state, path, transparent_bounce);
uint32_t shadow_flag = INTEGRATOR_STATE(state, path, flag);
const Spectrum throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
const Spectrum phase_sum = bsdf_eval_sum(&phase_eval);
const Spectrum throughput_phase = throughput * phase_sum;
if (!(kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_TREE)) {
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, bsdf_eval_average) = average(phase_sum);
}
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
PackedSpectrum pass_diffuse_weight;

6
intern/cycles/kernel/integrator/shadow_state_template.h
View file

@ -59,6 +59,12 @@ KERNEL_STRUCT_MEMBER(shadow_path,
KERNEL_STRUCT_MEMBER(shadow_path, uint64_t, path_segment, KERNEL_FEATURE_PATH_GUIDING)
#endif
KERNEL_STRUCT_MEMBER(shadow_path, float, guiding_mis_weight, KERNEL_FEATURE_PATH_GUIDING)
/* Only need when path tracing without the light tree. Stored as a single float to save
* space, as we do not expect to make it a big difference. */
KERNEL_STRUCT_MEMBER(shadow_path,
float,
bsdf_eval_average,
KernelFeatureRequest(KERNEL_FEATURE_PATH_TRACING, KERNEL_FEATURE_LIGHT_TREE))
KERNEL_STRUCT_END(shadow_path)
/********************************** Shadow Ray *******************************/

66
intern/cycles/kernel/light/area.h
View file

@ -247,12 +247,12 @@ ccl_device_forceinline bool area_light_is_ellipse(const ccl_global KernelAreaLig
/* Common API. */
/* Compute `eval_fac` and `pdf`. Also sample a new position on the light if `sample_coord`. */
template<bool in_volume_segment>
ccl_device_inline bool area_light_eval(const ccl_global KernelLight *klight,
const float3 ray_P,
ccl_private float3 *light_P,
ccl_private LightSample *ccl_restrict ls,
const float2 rand,
bool sample_coord)
ccl_device_forceinline bool area_light_eval(const ccl_global KernelLight *klight,
const float3 ray_P,
ccl_private float3 *light_P,
ccl_private LightSample *ccl_restrict ls,
const float2 rand,
bool sample_coord)
{
float3 axis_u = klight->area.axis_u;
float3 axis_v = klight->area.axis_v;
@ -368,10 +368,6 @@ ccl_device_inline bool area_light_sample(const ccl_global KernelLight *klight,
light_u /= klight->area.len_u;
light_v /= klight->area.len_v;
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
ls->u = light_v + 0.5f;
ls->v = -light_u - light_v;
return true;
}
@ -393,9 +389,7 @@ ccl_device_forceinline void area_light_mnee_sample_update(const ccl_global Kerne
ccl_device_inline bool area_light_intersect(const ccl_global KernelLight *klight,
const ccl_private Ray *ccl_restrict ray,
ccl_private float *t,
ccl_private float *u,
ccl_private float *v)
ccl_private float *t)
{
/* Area light. */
const float invarea = fabsf(klight->area.invarea);
@ -416,6 +410,7 @@ ccl_device_inline bool area_light_intersect(const ccl_global KernelLight *klight
const float3 light_P = klight->co;
float3 P;
float u, v;
return ray_quad_intersect(ray->P,
ray->D,
ray->tmin,
@ -426,25 +421,44 @@ ccl_device_inline bool area_light_intersect(const ccl_global KernelLight *klight
Ng,
&P,
t,
u,
v,
&u,
&v,
is_ellipse);
}
ccl_device_inline bool area_light_sample_from_intersection(
const ccl_global KernelLight *klight,
const ccl_private Intersection *ccl_restrict isect,
const float3 ray_P,
const float3 ray_D,
ccl_private LightSample *ccl_restrict ls)
ccl_device_inline float2 area_light_uv(const ccl_global KernelLight *klight, const float3 P)
{
ls->u = isect->u;
ls->v = isect->v;
ls->D = ray_D;
ls->Ng = klight->area.dir;
/* Compute uv when we already know there is an intersection, to avoid the need
* of storing this in the integrate state. */
const float3 inv_extent_u = klight->area.axis_u / klight->area.len_u;
const float3 inv_extent_v = klight->area.axis_v / klight->area.len_v;
const float3 light_P = klight->co;
const float3 inplane = P - light_P;
const float u = clamp(dot(inplane, inv_extent_u), -0.5f, 0.5f);
const float v = clamp(dot(inplane, inv_extent_v), -0.5f, 0.5f);
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
return make_float2(v + 0.5f, -u - v);
}
ccl_device_inline LightEval area_light_eval_from_intersection(const ccl_global KernelLight *klight,
const float3 ray_P,
const float3 ray_D,
const float t)
{
LightSample ls{};
ls.t = t;
ls.P = ray_P + ray_D * t;
ls.D = ray_D;
ls.Ng = klight->area.dir;
float3 light_P = klight->co;
return area_light_eval<false>(klight, ray_P, &light_P, ls, zero_float2(), false);
if (!area_light_eval<false>(klight, ray_P, &light_P, &ls, zero_float2(), false)) {
return LightEval{};
}
return LightEval{ls.eval_fac, ls.pdf};
}
/* Returns the maximal distance between the light center and the boundary. */

15
intern/cycles/kernel/light/common.h
View file

@ -10,17 +10,18 @@
CCL_NAMESPACE_BEGIN
/* Light Sample Result */
/* Result from light sampling with next event estimation.
*
* TODO: It may be possible to reduce the size of this struct now that shader evaluation
* no longer uses this. For example D, Ng or P, though it's not trivial. */
struct LightSample {
float3 P; /* position on light, or direction for distant light */
packed_float3 Ng; /* normal on light */
float t; /* distance to light (FLT_MAX for distant light) */
float3 D; /* direction from shading point to light */
float u, v; /* parametric coordinate on primitive */
float pdf; /* pdf for selecting light and point on light */
float pdf_selection; /* pdf for selecting light */
float eval_fac; /* intensity multiplier */
float eval_fac; /* intensity multiplier (normalization, spot falloff) */
int object; /* object id for triangle/curve lights */
int prim; /* lamp id for lights, primitive id for triangle/curve lights */
int shader; /* shader id */
@ -29,6 +30,12 @@ struct LightSample {
int emitter_id; /* index in the emitter array */
};
/* Result of evaluating a light from an intersection. */
struct LightEval {
float eval_fac = 0.0f; /* Intensity multiplier (normalization, spot falloff) */
float pdf = 0.0f; /* Pdf for light sampling with next event estimation sampling. */
};
/* Utilities */
ccl_device_inline float3 ellipse_sample(const float3 ru, const float3 rv, const float2 rand)

69
intern/cycles/kernel/light/distant.h
View file

@ -12,11 +12,9 @@
CCL_NAMESPACE_BEGIN
ccl_device_inline void distant_light_uv(KernelGlobals kg,
const ccl_global KernelLight *klight,
const float3 D,
ccl_private float *u,
ccl_private float *v)
ccl_device_inline float2 distant_light_uv(KernelGlobals kg,
const ccl_global KernelLight *klight,
const float3 D)
{
/* Map direction (x, y, z) to disk [-0.5, 0.5]^2:
* r^2 = (1 - z) / (1 - cos(klight->distant.angle))
@ -30,12 +28,10 @@ ccl_device_inline void distant_light_uv(KernelGlobals kg,
const float v_ = dot(D, make_float3(itfm.y)) * fac;
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
*u = v_ + 0.5f;
*v = -u_ - v_;
return make_float2(v_ + 0.5f, -u_ - v_);
}
ccl_device_inline bool distant_light_sample(KernelGlobals kg,
const ccl_global KernelLight *klight,
ccl_device_inline bool distant_light_sample(const ccl_global KernelLight *klight,
const float2 rand,
ccl_private LightSample *ls)
{
@ -49,22 +45,18 @@ ccl_device_inline bool distant_light_sample(KernelGlobals kg,
ls->eval_fac = klight->distant.eval_fac;
distant_light_uv(kg, klight, ls->D, &ls->u, &ls->v);
return true;
}
/* Special intersection check.
* Returns true if the distant_light_sample_from_intersection() for this light would return true.
* Returns true if the distant_light_eval_from_intersection() for this light would return true.
*
* The intersection parameters t, u, v are optimized for the shadow ray towards a dedicated light:
* u = v = 0, t = FLT_MAX.
*/
ccl_device bool distant_light_intersect(const ccl_global KernelLight *klight,
const ccl_private Ray *ccl_restrict ray,
ccl_private float *t,
ccl_private float *u,
ccl_private float *v)
ccl_private float *t)
{
kernel_assert(klight->type == LIGHT_DISTANT);
@ -77,59 +69,22 @@ ccl_device bool distant_light_intersect(const ccl_global KernelLight *klight,
}
*t = FLT_MAX;
*u = 0.0f;
*v = 0.0f;
return true;
}
ccl_device bool distant_light_sample_from_intersection(KernelGlobals kg,
const float3 ray_D,
const int lamp,
ccl_private LightSample *ccl_restrict ls)
ccl_device LightEval distant_light_eval_from_intersection(const ccl_global KernelLight *klight,
const float3 ray_D)
{
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, lamp);
const int shader = klight->shader_id;
const LightType type = (LightType)klight->type;
if (type != LIGHT_DISTANT) {
return false;
}
if (!(shader & SHADER_USE_MIS)) {
return false;
}
if (klight->distant.angle == 0.0f) {
return false;
return LightEval{};
}
/* Workaround to prevent a hang in the classroom scene with AMD HIP drivers 22.10,
* Remove when a compiler fix is available. */
#ifdef __HIP__
ls->shader = klight->shader_id;
#endif
if (vector_angle(-klight->co, ray_D) > klight->distant.angle) {
return false;
return LightEval{};
}
ls->type = type;
#ifndef __HIP__
ls->shader = klight->shader_id;
#endif
ls->object = klight->object_id;
ls->prim = lamp;
ls->t = FLT_MAX;
ls->P = -ray_D;
ls->Ng = -ray_D;
ls->D = ray_D;
ls->group = object_lightgroup(kg, ls->object);
ls->pdf = klight->distant.pdf;
ls->eval_fac = klight->distant.eval_fac;
distant_light_uv(kg, klight, ray_D, &ls->u, &ls->v);
return true;
return LightEval{klight->distant.eval_fac, klight->distant.pdf};
}
template<bool in_volume_segment>

94
intern/cycles/kernel/light/light.h
View file

@ -16,6 +16,7 @@
#include "kernel/light/spot.h"
#include "kernel/light/triangle.h"
#include "kernel/sample/lcg.h"
#include "kernel/types.h"
CCL_NAMESPACE_BEGIN
@ -125,8 +126,6 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
ls->shader = klight->shader_id;
ls->object = klight->object_id;
ls->prim = lamp;
ls->u = rand.x;
ls->v = rand.y;
ls->group = object_lightgroup(kg, ls->object);
if (in_volume_segment && (type == LIGHT_DISTANT || type == LIGHT_BACKGROUND)) {
@ -143,7 +142,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
}
if (type == LIGHT_DISTANT) {
if (!distant_light_sample(kg, klight, rand, ls)) {
if (!distant_light_sample(klight, rand, ls)) {
return false;
}
}
@ -163,7 +162,7 @@ ccl_device_inline bool light_sample(KernelGlobals kg,
}
}
else if (type == LIGHT_POINT) {
if (!point_light_sample(kg, klight, rand, P, N, shader_flags, ls)) {
if (!point_light_sample(klight, rand, P, N, shader_flags, ls)) {
return false;
}
}
@ -329,8 +328,6 @@ ccl_device_forceinline int lights_intersect_impl(KernelGlobals kg,
const LightType type = (LightType)klight->type;
float t = 0.0f;
float u = 0.0f;
float v = 0.0f;
if (type == LIGHT_SPOT) {
if (!spot_light_intersect(klight, ray, &t)) {
@ -343,7 +340,7 @@ ccl_device_forceinline int lights_intersect_impl(KernelGlobals kg,
}
}
else if (type == LIGHT_AREA) {
if (!area_light_intersect(klight, ray, &t, &u, &v)) {
if (!area_light_intersect(klight, ray, &t)) {
continue;
}
}
@ -351,7 +348,7 @@ ccl_device_forceinline int lights_intersect_impl(KernelGlobals kg,
if (is_main_path || ray->tmax != FLT_MAX) {
continue;
}
if (!distant_light_intersect(klight, ray, &t, &u, &v)) {
if (!distant_light_intersect(klight, ray, &t)) {
continue;
}
}
@ -384,8 +381,8 @@ ccl_device_forceinline int lights_intersect_impl(KernelGlobals kg,
}
isect->t = t;
isect->u = u;
isect->v = v;
isect->u = 0.0f;
isect->v = 0.0f;
isect->type = PRIMITIVE_LAMP;
isect->prim = lamp;
isect->object = object;
@ -454,46 +451,61 @@ ccl_device int lights_intersect_shadow_linked(KernelGlobals kg,
/* Setup light sample from intersection. */
ccl_device bool light_sample_from_intersection(KernelGlobals kg,
const ccl_private Intersection *ccl_restrict isect,
const float3 ray_P,
const float3 ray_D,
const float3 N,
const uint32_t path_flag,
ccl_private LightSample *ccl_restrict ls)
ccl_device LightEval
light_eval_from_intersection(KernelGlobals kg,
const ccl_private Intersection *ccl_restrict isect,
const float3 ray_P,
const float3 ray_D,
const float3 N,
const uint32_t path_flag)
{
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, isect->prim);
const LightType type = (LightType)klight->type;
ls->type = type;
ls->shader = klight->shader_id;
ls->object = isect->object;
ls->prim = isect->prim;
ls->t = isect->t;
ls->P = ray_P + ray_D * ls->t;
ls->D = ray_D;
ls->group = object_lightgroup(kg, ls->object);
if (type == LIGHT_SPOT) {
if (!spot_light_sample_from_intersection(kg, klight, ray_P, ray_D, N, path_flag, ls)) {
return false;
}
return spot_light_eval_from_intersection(kg, klight, ray_P, ray_D, isect->t, N, path_flag);
}
else if (type == LIGHT_POINT) {
if (!point_light_sample_from_intersection(kg, klight, ray_P, ray_D, N, path_flag, ls)) {
return false;
}
if (type == LIGHT_POINT) {
return point_light_eval_from_intersection(klight, ray_P, ray_D, isect->t, N, path_flag);
}
else if (type == LIGHT_AREA) {
if (!area_light_sample_from_intersection(klight, isect, ray_P, ray_D, ls)) {
return false;
}
}
else {
kernel_assert(!"Invalid lamp type in light_sample_from_intersection");
return false;
if (type == LIGHT_AREA) {
return area_light_eval_from_intersection(klight, ray_P, ray_D, isect->t);
}
return true;
kernel_assert(!"Invalid lamp type in light_eval_from_intersection");
return LightEval{};
}
/* Get light coordinates from position on light. */
ccl_device void light_normal_uv_from_position(KernelGlobals kg,
const ccl_global KernelLight *klight,
const float3 P,
const float3 D,
ccl_private float3 &Ng,
ccl_private float2 &uv)
{
const LightType type = (LightType)klight->type;
if (type == LIGHT_SPOT) {
Ng = (klight->spot.is_sphere) ? normalize(P - klight->co) : -D;
const float3 local_ray = spot_light_to_local(kg, klight, -D);
uv = spot_light_uv(local_ray, klight->spot.half_cot_half_spot_angle);
}
else if (type == LIGHT_POINT) {
Ng = (klight->spot.is_sphere) ? normalize(P - klight->co) : -D;
uv = point_light_uv(kg, klight, Ng);
}
else if (type == LIGHT_AREA) {
Ng = klight->area.dir;
uv = area_light_uv(klight, P);
}
else if (type == LIGHT_DISTANT) {
Ng = -D;
uv = distant_light_uv(kg, klight, D);
}
else {
kernel_assert(0);
}
}
CCL_NAMESPACE_END

69
intern/cycles/kernel/light/point.h
View file

@ -10,12 +10,12 @@
#include "kernel/light/common.h"
#include "util/defines.h"
#include "util/math_intersect.h"
CCL_NAMESPACE_BEGIN
ccl_device_inline bool point_light_sample(KernelGlobals kg,
const ccl_global KernelLight *klight,
ccl_device_inline bool point_light_sample(const ccl_global KernelLight *klight,
const float2 rand,
const float3 P,
const float3 N,
@ -76,13 +76,6 @@ ccl_device_inline bool point_light_sample(KernelGlobals kg,
ls->pdf = invarea * light_pdf_area_to_solid_angle(lightN, -ls->D, ls->t);
}
/* Texture coordinates. */
const Transform itfm = lamp_get_inverse_transform(kg, klight);
const float2 uv = map_to_sphere(transform_direction(&itfm, ls->Ng));
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
ls->u = uv.y;
ls->v = 1.0f - uv.x - uv.y;
return true;
}
@ -97,8 +90,18 @@ ccl_device_forceinline float sphere_light_pdf(
return has_transmission ? M_1_2PI_F * 0.5f : pdf_cos_hemisphere(N, D);
}
ccl_device_forceinline void point_light_mnee_sample_update(KernelGlobals kg,
const ccl_global KernelLight *klight,
ccl_device_forceinline float2 point_light_uv(KernelGlobals kg,
const ccl_global KernelLight *klight,
const float3 Ng)
{
/* Texture coordinates. */
const Transform itfm = lamp_get_inverse_transform(kg, klight);
const float2 uv = map_to_sphere(transform_direction(&itfm, Ng));
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
return make_float2(uv.y, 1.0f - uv.x - uv.y);
}
ccl_device_forceinline void point_light_mnee_sample_update(const ccl_global KernelLight *klight,
ccl_private LightSample *ls,
const float3 P,
const float3 N,
@ -126,13 +129,6 @@ ccl_device_forceinline void point_light_mnee_sample_update(KernelGlobals kg,
ls->Ng = -ls->D;
}
/* Texture coordinates. */
const Transform itfm = lamp_get_inverse_transform(kg, klight);
const float2 uv = map_to_sphere(transform_direction(&itfm, ls->Ng));
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
ls->u = uv.y;
ls->v = 1.0f - uv.x - uv.y;
}
ccl_device_inline bool point_light_intersect(const ccl_global KernelLight *klight,
@ -155,44 +151,31 @@ ccl_device_inline bool point_light_intersect(const ccl_global KernelLight *kligh
ray->P, ray->D, ray->tmin, ray->tmax, klight->co, diskN, radius, &P, t);
}
ccl_device_inline bool point_light_sample_from_intersection(KernelGlobals kg,
const ccl_global KernelLight *klight,
const float3 ray_P,
const float3 ray_D,
const float3 N,
const uint32_t path_flag,
ccl_private LightSample *ccl_restrict
ls)
ccl_device_inline LightEval
point_light_eval_from_intersection(const ccl_global KernelLight *klight,
const float3 ray_P,
const float3 ray_D,
const float t,
const float3 N,
const uint32_t path_flag)
{
const float r_sq = sqr(klight->spot.radius);
ls->eval_fac = klight->spot.eval_fac;
LightEval light_eval = {klight->spot.eval_fac, 0.0f};
if (klight->spot.is_sphere) {
const float d_sq = len_squared(ray_P - klight->co);
ls->pdf = sphere_light_pdf(d_sq, r_sq, N, ray_D, path_flag);
ls->Ng = normalize(ls->P - klight->co);
light_eval.pdf = sphere_light_pdf(d_sq, r_sq, N, ray_D, path_flag);
}
else {
if (ls->t != FLT_MAX) {
if (t != FLT_MAX) {
const float3 lightN = normalize(ray_P - klight->co);
const float invarea = (r_sq > 0.0f) ? 1.0f / (r_sq * M_PI_F) : 1.0f;
ls->pdf = invarea * light_pdf_area_to_solid_angle(lightN, -ray_D, ls->t);
light_eval.pdf = invarea * light_pdf_area_to_solid_angle(lightN, -ray_D, t);
}
else {
ls->pdf = 0.0f;
}
ls->Ng = -ray_D;
}
/* Texture coordinates. */
const Transform itfm = lamp_get_inverse_transform(kg, klight);
const float2 uv = map_to_sphere(transform_direction(&itfm, ls->Ng));
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
ls->u = uv.y;
ls->v = 1.0f - uv.x - uv.y;
return true;
return light_eval;
}
template<bool in_volume_segment>

158
intern/cycles/kernel/light/sample.h
View file

@ -20,47 +20,67 @@
CCL_NAMESPACE_BEGIN
/* Evaluate shader on light. */
ccl_device_noinline_cpu Spectrum
light_sample_shader_eval(KernelGlobals kg,
IntegratorState state,
ccl_private ShaderData *ccl_restrict emission_sd,
ccl_private LightSample *ccl_restrict ls,
const float time)
/* Evaluate constant factors for a direct light sample. */
ccl_device bool light_sample_shader_eval_nee_constant(KernelGlobals kg,
const int shader_id,
const int prim,
const bool is_light,
ccl_private Spectrum &eval)
{
eval = one_spectrum();
const bool is_constant = surface_shader_constant_emission(kg, shader_id, &eval);
if (is_light) {
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, prim);
eval *= rgb_to_spectrum(
make_float3(klight->strength[0], klight->strength[1], klight->strength[2]));
}
return is_constant;
}
/* Evaluate shader on light. Not supported for background and triangle lights, that happens
* in shade_surface and shader_background. */
ccl_device_noinline_cpu Spectrum light_sample_shader_eval_forward(KernelGlobals kg,
IntegratorState state,
const int light_id,
const float3 ray_P,
const float3 ray_D,
const float t,
const float time)
{
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, light_id);
/* setup shading at emitter */
Spectrum eval = zero_spectrum();
if (surface_shader_constant_emission(kg, ls->shader, &eval)) {
if ((ls->prim != PRIM_NONE) && dot(ls->Ng, ls->D) > 0.0f) {
ls->Ng = -ls->Ng;
}
}
else {
if (!surface_shader_constant_emission(kg, klight->shader_id, &eval)) {
/* Setup shader data and call surface_shader_eval once, better
* for GPU coherence and compile times. */
PROFILING_INIT_FOR_SHADER(kg, PROFILING_SHADE_LIGHT_SETUP);
if (ls->type == LIGHT_BACKGROUND) {
shader_setup_from_background(kg, emission_sd, ls->P, ls->D, time);
}
else {
shader_setup_from_sample(kg,
emission_sd,
ls->P,
ls->Ng,
-ls->D,
ls->shader,
ls->object,
ls->prim,
ls->u,
ls->v,
ls->t,
time,
false,
ls->type != LIGHT_TRIANGLE);
ls->Ng = emission_sd->Ng;
}
ShaderDataTinyStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
const float3 P = (t == FLT_MAX) ? -ray_D : ray_P + ray_D * t;
float3 Ng = zero_float3();
float2 uv = zero_float2();
light_normal_uv_from_position(kg, klight, P, ray_D, Ng, uv);
shader_setup_from_sample(kg,
emission_sd,
P,
Ng,
-ray_D,
klight->shader_id,
klight->object_id,
light_id,
uv.x,
uv.y,
t,
time,
false,
true);
PROFILING_SHADER(emission_sd->object, emission_sd->shader);
PROFILING_EVENT(PROFILING_SHADE_LIGHT_EVAL);
@ -71,18 +91,11 @@ light_sample_shader_eval(KernelGlobals kg,
kg, state, emission_sd, nullptr, PATH_RAY_EMISSION);
/* Evaluate closures. */
if (ls->type == LIGHT_BACKGROUND) {
eval = surface_shader_background(emission_sd);
}
else {
eval = surface_shader_emission(emission_sd);
}
eval = surface_shader_emission(emission_sd);
}
eval *= ls->eval_fac;
if (ls->type != LIGHT_TRIANGLE) {
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, ls->prim);
{
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, light_id);
eval *= rgb_to_spectrum(
make_float3(klight->strength[0], klight->strength[1], klight->strength[2]));
}
@ -91,6 +104,14 @@ light_sample_shader_eval(KernelGlobals kg,
}
/* Early path termination of shadow rays. */
ccl_device_inline float light_sample_terminate_probability(KernelGlobals kg,
ccl_private Spectrum eval)
{
return (kernel_data.integrator.light_inv_rr_threshold > 0.0f) ?
reduce_max(fabs(eval)) * kernel_data.integrator.light_inv_rr_threshold :
1.0f;
}
ccl_device_inline bool light_sample_terminate(KernelGlobals kg,
ccl_private BsdfEval *ccl_restrict eval,
const float rand_terminate)
@ -99,15 +120,35 @@ ccl_device_inline bool light_sample_terminate(KernelGlobals kg,
return true;
}
if (kernel_data.integrator.light_inv_rr_threshold > 0.0f) {
const float probability = reduce_max(fabs(bsdf_eval_sum(eval))) *
kernel_data.integrator.light_inv_rr_threshold;
if (probability < 1.0f) {
if (rand_terminate >= probability) {
return true;
}
bsdf_eval_mul(eval, 1.0f / probability);
const float probability = light_sample_terminate_probability(kg, bsdf_eval_sum(eval));
if (probability < 1.0f) {
if (rand_terminate >= probability) {
return true;
}
bsdf_eval_mul(eval, 1.0f / probability);
}
return false;
}
ccl_device_inline bool light_sample_terminate(KernelGlobals kg,
ccl_private Spectrum &light_eval,
const float bsdf_eval,
const float rand_terminate)
{
/* Same logic as above, but where bsdf_eval is already part of the throughput so
* we only need to modify the light eval while still taking into account bsdf eval
* for the termination probability. */
if (is_zero(light_eval)) {
return true;
}
const float probability = light_sample_terminate_probability(kg, light_eval * bsdf_eval);
if (probability < 1.0f) {
if (rand_terminate >= probability) {
return true;
}
light_eval /= probability;
}
return false;
@ -396,7 +437,7 @@ ccl_device_forceinline void light_sample_update(KernelGlobals kg,
const ccl_global KernelLight *klight = &kernel_data_fetch(lights, ls->prim);
if (ls->type == LIGHT_POINT) {
point_light_mnee_sample_update(kg, klight, ls, P, N, path_flag);
point_light_mnee_sample_update(klight, ls, P, N, path_flag);
}
else if (ls->type == LIGHT_SPOT) {
spot_light_mnee_sample_update(kg, klight, ls, P, N, path_flag);
@ -467,7 +508,8 @@ ccl_device_inline float light_sample_mis_weight_forward_surface(KernelGlobals kg
ccl_device_inline float light_sample_mis_weight_forward_lamp(KernelGlobals kg,
IntegratorState state,
const uint32_t path_flag,
const ccl_private LightSample *ls,
const int light_id,
const float light_sample_pdf,
const float3 P)
{
if (path_flag & PATH_RAY_MIS_SKIP) {
@ -475,7 +517,7 @@ ccl_device_inline float light_sample_mis_weight_forward_lamp(KernelGlobals kg,
}
const float mis_ray_pdf = INTEGRATOR_STATE(state, path, mis_ray_pdf);
float pdf = ls->pdf;
float pdf = light_sample_pdf;
/* Light selection pdf. */
#ifdef __LIGHT_TREE__
@ -488,7 +530,7 @@ ccl_device_inline float light_sample_mis_weight_forward_lamp(KernelGlobals kg,
dt,
path_flag,
0,
kernel_data_fetch(light_to_tree, ls->prim),
kernel_data_fetch(light_to_tree, light_id),
light_link_receiver_forward(kg, state));
}
else
@ -503,10 +545,12 @@ ccl_device_inline float light_sample_mis_weight_forward_lamp(KernelGlobals kg,
ccl_device_inline float light_sample_mis_weight_forward_distant(KernelGlobals kg,
IntegratorState state,
const uint32_t path_flag,
const ccl_private LightSample *ls)
const int light_id,
const float light_sample_pdf)
{
const float3 ray_P = INTEGRATOR_STATE(state, ray, P);
return light_sample_mis_weight_forward_lamp(kg, state, path_flag, ls, ray_P);
return light_sample_mis_weight_forward_lamp(
kg, state, path_flag, light_id, light_sample_pdf, ray_P);
}
ccl_device_inline float light_sample_mis_weight_forward_background(KernelGlobals kg,

55
intern/cycles/kernel/light/spot.h
View file

@ -30,17 +30,13 @@ ccl_device float spot_light_attenuation(const ccl_global KernelSpotLight *spot,
return smoothstepf((ray.z - spot->cos_half_spot_angle) * spot->spot_smooth);
}
ccl_device void spot_light_uv(const float3 ray,
const float half_cot_half_spot_angle,
ccl_private float *u,
ccl_private float *v)
ccl_device float2 spot_light_uv(const float3 ray, const float half_cot_half_spot_angle)
{
/* Ensures that the spot light projects the full image regardless of the spot angle. */
const float factor = half_cot_half_spot_angle / ray.z;
/* NOTE: Return barycentric coordinates in the same notation as Embree and OptiX. */
*u = ray.y * factor + 0.5f;
*v = -(ray.x + ray.y) * factor;
return make_float2(ray.y * factor + 0.5f, -(ray.x + ray.y) * factor);
}
template<bool in_volume_segment>
@ -122,9 +118,6 @@ ccl_device_inline bool spot_light_sample(KernelGlobals kg,
/* Remap sampled point onto the sphere to prevent precision issues with small radius. */
ls->Ng = normalize(ls->P - klight->co);
ls->P = ls->Ng * klight->spot.radius + klight->co;
/* Texture coordinates. */
spot_light_uv(local_ray, klight->spot.half_cot_half_spot_angle, &ls->u, &ls->v);
}
else {
/* Point light with ad-hoc radius based on oriented disk. */
@ -146,9 +139,6 @@ ccl_device_inline bool spot_light_sample(KernelGlobals kg,
/* PDF. */
const float invarea = (r_sq > 0.0f) ? 1.0f / (r_sq * M_PI_F) : 1.0f;
ls->pdf = invarea * light_pdf_area_to_solid_angle(lightN, -ls->D, ls->t);
/* Texture coordinates. */
spot_light_uv(local_ray, klight->spot.half_cot_half_spot_angle, &ls->u, &ls->v);
}
return true;
@ -211,9 +201,6 @@ ccl_device_forceinline void spot_light_mnee_sample_update(KernelGlobals kg,
if (use_attenuation) {
ls->eval_fac *= spot_light_attenuation(&klight->spot, local_ray);
}
/* Texture coordinates. */
spot_light_uv(local_ray, klight->spot.half_cot_half_spot_angle, &ls->u, &ls->v);
}
ccl_device_inline bool spot_light_intersect(const ccl_global KernelLight *klight,
@ -228,49 +215,37 @@ ccl_device_inline bool spot_light_intersect(const ccl_global KernelLight *klight
return point_light_intersect(klight, ray, t);
}
ccl_device_inline bool spot_light_sample_from_intersection(KernelGlobals kg,
const ccl_global KernelLight *klight,
const float3 ray_P,
const float3 ray_D,
const float3 N,
const uint32_t path_flag,
ccl_private LightSample *ccl_restrict
ls)
ccl_device_inline LightEval spot_light_eval_from_intersection(KernelGlobals kg,
const ccl_global KernelLight *klight,
const float3 ray_P,
const float3 ray_D,
const float t,
const float3 N,
const uint32_t path_flag)
{
const float r_sq = sqr(klight->spot.radius);
const float d_sq = len_squared(ray_P - klight->co);
ls->eval_fac = klight->spot.eval_fac;
LightEval light_eval = {klight->spot.eval_fac, 0.0f};
if (klight->spot.is_sphere) {
ls->pdf = spot_light_pdf(&klight->spot, d_sq, r_sq, N, ray_D, path_flag);
ls->Ng = normalize(ls->P - klight->co);
light_eval.pdf = spot_light_pdf(&klight->spot, d_sq, r_sq, N, ray_D, path_flag);
}
else {
if (ls->t != FLT_MAX) {
if (t != FLT_MAX) {
const float3 lightN = normalize(ray_P - klight->co);
const float invarea = (r_sq > 0.0f) ? 1.0f / (r_sq * M_PI_F) : 1.0f;
ls->pdf = invarea * light_pdf_area_to_solid_angle(lightN, -ray_D, ls->t);
light_eval.pdf = invarea * light_pdf_area_to_solid_angle(lightN, -ray_D, t);
}
else {
ls->pdf = 0.0f;
}
ls->Ng = -ray_D;
}
/* Attenuation. */
const float3 local_ray = spot_light_to_local(kg, klight, -ray_D);
if (!klight->spot.is_sphere || d_sq > r_sq) {
ls->eval_fac *= spot_light_attenuation(&klight->spot, local_ray);
}
if (ls->eval_fac == 0) {
return false;
light_eval.eval_fac *= spot_light_attenuation(&klight->spot, local_ray);
}
/* Texture coordinates. */
spot_light_uv(local_ray, klight->spot.half_cot_half_spot_angle, &ls->u, &ls->v);
return true;
return light_eval;
}
/* Find the ray segment lit by the spot light. */

18
intern/cycles/kernel/light/triangle.h
View file

@ -217,7 +217,9 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
ls->D = z * B + sin_from_cos(z) * safe_normalize(C_ - dot(C_, B) * B);
/* calculate intersection with the planar triangle */
if (!ray_triangle_intersect(P, ls->D, 0.0f, FLT_MAX, V[0], V[1], V[2], &ls->u, &ls->v, &ls->t))
float unused_u, unused_v;
if (!ray_triangle_intersect(
P, ls->D, 0.0f, FLT_MAX, V[0], V[1], V[2], &unused_u, &unused_v, &ls->t))
{
ls->pdf = 0.0f;
return false;
@ -256,8 +258,6 @@ ccl_device_forceinline bool triangle_light_sample(KernelGlobals kg,
/* compute incoming direction, distance and pdf */
ls->D = normalize_len(ls->P - P, &ls->t);
ls->pdf = triangle_light_pdf_area_sampling(ls->Ng, -ls->D, ls->t) / area;
ls->u = u;
ls->v = v;
}
/* Belongs in distribution.h but can reuse computations here. */
@ -336,4 +336,16 @@ ccl_device_forceinline bool triangle_light_tree_parameters(
return front_facing && shape_above_surface;
}
ccl_device float2 triangle_light_uv(KernelGlobals kg,
const int object,
const int prim,
const float time,
const float3 ray_P,
const float3 ray_D)
{
float3 V[3];
triangle_world_space_vertices(kg, object, prim, time, V);
return ray_triangle_uv(ray_P, ray_D, V[0], V[1], V[2]);
}
CCL_NAMESPACE_END

4
intern/cycles/kernel/sample/mapping.h
View file

@ -10,6 +10,10 @@
#include "util/math.h"
#include "util/projection.h"
#ifndef __KERNEL_GPU__
# include <climits>
#endif
CCL_NAMESPACE_BEGIN
/* Distribute 2D uniform random samples on [0, 1] over unit disk [-1, 1], with concentric mapping

3
intern/cycles/kernel/types.h
View file

@ -1646,7 +1646,8 @@ enum DeviceKernel : int {
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_LIGHT_NEE,
DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE,

28
intern/cycles/util/math_intersect.h
View file

@ -221,6 +221,34 @@ ccl_device_forceinline bool ray_triangle_intersect(const float3 ray_P,
return true;
}
ccl_device_forceinline float2 ray_triangle_uv(const float3 ray_P,
const float3 ray_D,
const float3 tri_a,
const float3 tri_b,
const float3 tri_c)
{
/* Matches ray_triangle-intersect, for an intersection known to exist. */
/* Calculate vertices relative to ray origin. */
const float3 v0 = tri_a - ray_P;
const float3 v1 = tri_b - ray_P;
const float3 v2 = tri_c - ray_P;
/* Calculate triangle edges. */
const float3 e0 = v2 - v0;
const float3 e1 = v0 - v1;
const float3 e2 = v1 - v2;
/* Compute barycentric coordinates. */
const float U = ray_triangle_dot(ray_triangle_cross(e0, v2 + v0), ray_D);
const float V = ray_triangle_dot(ray_triangle_cross(e1, v0 + v1), ray_D);
const float W = ray_triangle_dot(ray_triangle_cross(e2, v1 + v2), ray_D);
const float UVW = U + V + W;
const float rcp_uvw = (fabsf(UVW) < 1e-18f) ? 0.0f : ray_triangle_reciprocal(UVW);
return make_float2(clamp(U * rcp_uvw, 0.0f, 1.0f), clamp(V * rcp_uvw, 0.0f, 1.0f));
}
ccl_device_forceinline bool ray_triangle_intersect_self(const float3 ray_P,
const float3 ray_D,
const float3 verts[3])