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

Refactor: Cycles: Move MNEE walk into separate kernel

Browse source 
The new intersect_mnee kernel runs before shade_surface, and
shade_surface_mnee is eliminated. That large kernel was causing problems
for some GPU compilers.

MNEE state is packed into a shadow path state to avoid significantly
increasing the path state size. This shadow state is then either turned
into an actual shadow ray state or discarded in shade_surface.

MNEE was re-enabled on HIP RDNA2 as it works again now. Texture cache
misses now also work correctly with MNEE.

This adds some extra code to the regular shade_surface kernel even when
MNEE is not used, to use the MNEE sampled point instead of sampling a
light. But there seems to be no significant performance impact.

Co-authored-by: Sergey Sharybin <sergey@blender.org>
Pull Request: https://projects.blender.org/blender/blender/pulls/158698
This commit is contained in:
Brecht Van Lommel 2026-05-27 20:34:06 +02:00
parent 428c9987f5
commit 5fc85d1cb7
34 changed files with 487 additions and 205 deletions
Download patch file Download diff file Expand all files Collapse all files

2
intern/cycles/device/cuda/device_impl.cpp
View file

@ -484,8 +484,6 @@ void CUDADevice::reserve_local_memory(const uint kernel_features)
/* Use the biggest kernel for estimation. */
const DeviceKernel test_kernel = (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE :
(kernel_features & KERNEL_FEATURE_MNEE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE :
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE;
/* Launch kernel, using just 1 block appears sufficient to reserve memory for all

3
intern/cycles/device/cuda/kernel.cpp
View file

@ -16,8 +16,7 @@ void CUDADeviceKernels::load(CUDADevice *device)
for (int i = 0; i < (int)DEVICE_KERNEL_NUM; i++) {
CUDADeviceKernel &kernel = kernels_[i];
/* No mega-kernel used for GPU. */
if (i == DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL) {
if (!device_kernel_has_gpu_function((DeviceKernel)i)) {
continue;
}

4
intern/cycles/device/hip/device.cpp
View file

@ -167,11 +167,7 @@ void device_hip_info(vector<DeviceInfo> &devices)
info.description = string(name);
info.num = num;
# ifdef _WIN32
info.has_mnee = hipIsNotRDNA2(num);
# else
info.has_mnee = true;
# endif
info.has_nanovdb = true;
info.has_gpu_queue = true;

2
intern/cycles/device/hip/device_impl.cpp
View file

@ -431,8 +431,6 @@ void HIPDevice::reserve_local_memory(const uint kernel_features)
/* Use the biggest kernel for estimation. */
const DeviceKernel test_kernel = (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE :
(kernel_features & KERNEL_FEATURE_MNEE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE :
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE;
/* Launch kernel, using just 1 block appears sufficient to reserve memory for all

5
intern/cycles/device/hip/kernel.cpp
View file

@ -18,8 +18,7 @@ bool HIPDeviceKernels::load_kernel(HIPDevice *device,
return true;
}
/* No mega-kernel used for GPU. */
if (kernel == DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL) {
if (!device_kernel_has_gpu_function(kernel)) {
return false;
}
@ -62,9 +61,9 @@ void HIPDeviceKernels::load_raytrace(HIPDevice *device, hipModule_t hip_module)
load_kernel(device, hip_module, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE);
load_kernel(device, hip_module, DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK);
load_kernel(device, hip_module, DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT);
load_kernel(device, hip_module, DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
load_kernel(device, hip_module, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE);
load_kernel(device, hip_module, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE);
}
const HIPDeviceKernel &HIPDeviceKernels::get(DeviceKernel kernel) const

9
intern/cycles/device/hip/util.h
View file

@ -77,15 +77,6 @@ static inline bool hipIsRDNA2OrNewer(const int hipDevId)
return (major > 10 || (major == 10 && minor >= 3));
}
static inline bool hipIsNotRDNA2(const int hipDevId)
{
int major, minor;
hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, hipDevId);
hipDeviceGetAttribute(&minor, hipDeviceAttributeComputeCapabilityMinor, hipDevId);
return !(major == 10 && minor == 3);
}
static inline bool hipSupportsDeviceOIDN(const int hipDevId)
{
/* Matches HIPDevice::getArch in HIP. */

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

@ -17,7 +17,6 @@ bool device_kernel_has_shading(DeviceKernel kernel)
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 ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME_RAY_MARCHING ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW ||
@ -35,8 +34,14 @@ bool device_kernel_has_intersection(DeviceKernel kernel)
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);
kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE);
}
bool device_kernel_has_gpu_function(DeviceKernel kernel)
{
return !(kernel == DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING);
}
const char *device_kernel_as_string(DeviceKernel kernel)
@ -57,6 +62,10 @@ const char *device_kernel_as_string(DeviceKernel kernel)
return "integrator_intersect_volume_stack";
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
return "integrator_intersect_dedicated_light";
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE:
return "integrator_intersect_mnee";
case DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING:
return "integrator_shadow_path_mnee_pending";
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND:
return "integrator_shade_background";
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE:
@ -69,8 +78,6 @@ const char *device_kernel_as_string(DeviceKernel kernel)
return "integrator_shade_surface";
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
return "integrator_shade_surface_raytrace";
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
return "integrator_shade_surface_mnee";
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
return "integrator_shade_volume";
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME_RAY_MARCHING:

1
intern/cycles/device/kernel.h
View file

@ -19,6 +19,7 @@ CCL_NAMESPACE_BEGIN
bool device_kernel_has_shading(DeviceKernel kernel);
bool device_kernel_has_intersection(DeviceKernel kernel);
bool device_kernel_has_gpu_function(DeviceKernel kernel);
const char *device_kernel_as_string(DeviceKernel kernel);

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

@ -250,8 +250,8 @@ bool ShaderCache::should_load_kernel(DeviceKernel device_kernel,
return false;
}
if (device_kernel == DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL) {
/* Skip megakernel. */
if (!device_kernel_has_gpu_function(device_kernel)) {
/* Skip megakernel and other markers without a GPU function. */
return false;
}
@ -262,9 +262,9 @@ bool ShaderCache::should_load_kernel(DeviceKernel device_kernel,
}
}
if (device_kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE) {
if (device_kernel == DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE) {
if ((device->kernel_features & KERNEL_FEATURE_MNEE) == 0) {
/* Skip shade_surface_mnee kernel if the scene doesn't require it. */
/* Skip the MNEE kernel if the scene doesn't require it. */
return false;
}
}

4
intern/cycles/device/metal/queue.mm
View file

@ -634,7 +634,7 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
if (IntegratorQueueCounter *queue_counter = (IntegratorQueueCounter *)
it.first->host_pointer)
{
for (int i = 0; i < DEVICE_KERNEL_INTEGRATOR_NUM; i++) {
for (int i = 0; i < DEVICE_GPU_KERNEL_INTEGRATOR_NUM; i++) {
printf("%s%d", i == 0 ? "" : ",", queue_counter->num_queued[i]);
}
}
@ -816,7 +816,7 @@ void MetalDeviceQueue::prepare_resources(DeviceKernel /*kernel*/)
id<MTLComputeCommandEncoder> MetalDeviceQueue::get_compute_encoder(DeviceKernel kernel)
{
bool concurrent = int(kernel) < int(DEVICE_KERNEL_INTEGRATOR_NUM);
bool concurrent = int(kernel) < int(DEVICE_GPU_KERNEL_INTEGRATOR_NUM);
if (profiling_enabled_) {
/* Close the current encoder to ensure we're able to capture per-encoder timing data. */

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

@ -297,8 +297,6 @@ void OneapiDevice::reserve_private_memory(const uint kernel_features)
/* Use the biggest kernel for estimation. */
const DeviceKernel test_kernel = (kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE :
(kernel_features & KERNEL_FEATURE_MNEE) ?
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE :
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE;
{
@ -1316,6 +1314,7 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE:
preferred_work_group_size = preferred_work_group_size_intersect;
break;
@ -1324,7 +1323,6 @@ void OneapiDevice::get_adjusted_global_and_local_sizes(SyclQueue *queue,
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:
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME_RAY_MARCHING:
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:

4
intern/cycles/device/oneapi/graphics_interop.cpp
View file

@ -59,13 +59,13 @@ void OneapiDeviceGraphicsInterop::set_buffer(GraphicsInteropBuffer &interop_buff
sycl::ext::oneapi::experimental::external_mem_handle_type::win32_nt_handle;
sycl::ext::oneapi::experimental::external_mem_descriptor<
sycl::ext::oneapi::experimental::resource_win32_handle>
sycl_external_mem_descriptor{vulkan_windows_handle_, sycl_mem_handle_type};
sycl_external_mem_descriptor{{vulkan_windows_handle_}, sycl_mem_handle_type};
# else
/* import_external_memory will take ownership of the file descriptor. */
auto sycl_mem_handle_type = sycl::ext::oneapi::experimental::external_mem_handle_type::opaque_fd;
sycl::ext::oneapi::experimental::external_mem_descriptor<
sycl::ext::oneapi::experimental::resource_fd>
sycl_external_mem_descriptor{static_cast<int>(interop_buffer.take_handle()),
sycl_external_mem_descriptor{{static_cast<int>(interop_buffer.take_handle())},
sycl_mem_handle_type};
# endif

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

@ -547,10 +547,10 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
}
if (kernel_features & KERNEL_FEATURE_MNEE) {
group_descs[PG_RGEN_SHADE_SURFACE_MNEE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_SHADE_SURFACE_MNEE].raygen.module = optix_module;
group_descs[PG_RGEN_SHADE_SURFACE_MNEE].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_shade_surface_mnee";
group_descs[PG_RGEN_INTERSECT_MNEE].kind = OPTIX_PROGRAM_GROUP_KIND_RAYGEN;
group_descs[PG_RGEN_INTERSECT_MNEE].raygen.module = optix_module;
group_descs[PG_RGEN_INTERSECT_MNEE].raygen.entryFunctionName =
"__raygen__kernel_optix_integrator_intersect_mnee";
}
/* OSL uses direct callables to execute, so shading needs to be done in OptiX if OSL is used. */
@ -705,7 +705,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
pipeline_groups.push_back(groups[PG_CALL_SVM_BEVEL]);
}
if (kernel_features & KERNEL_FEATURE_MNEE) {
pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_MNEE]);
pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_MNEE]);
}
pipeline_groups.push_back(groups[PG_MISS]);
pipeline_groups.push_back(groups[PG_HITD]);
@ -757,7 +757,7 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
/* Combine ray generation and trace continuation stack size. */
const unsigned int css = std::max(stack_size[PG_RGEN_SHADE_SURFACE_RAYTRACE].cssRG,
stack_size[PG_RGEN_SHADE_SURFACE_MNEE].cssRG) +
stack_size[PG_RGEN_INTERSECT_MNEE].cssRG) +
link_options.maxTraceDepth * trace_css;
const unsigned int dss = std::max(stack_size[PG_CALL_SVM_AO].dssDC,
stack_size[PG_CALL_SVM_BEVEL].dssDC);
@ -1063,7 +1063,7 @@ bool OptiXDevice::load_osl_kernels()
pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_RAYTRACE]);
pipeline_groups.push_back(groups[PG_CALL_SVM_AO]);
pipeline_groups.push_back(groups[PG_CALL_SVM_BEVEL]);
pipeline_groups.push_back(groups[PG_RGEN_SHADE_SURFACE_MNEE]);
pipeline_groups.push_back(groups[PG_RGEN_INTERSECT_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]);
@ -1103,7 +1103,7 @@ bool OptiXDevice::load_osl_kernels()
}
const unsigned int css = std::max(stack_size[PG_RGEN_SHADE_SURFACE_RAYTRACE].cssRG,
stack_size[PG_RGEN_SHADE_SURFACE_MNEE].cssRG);
stack_size[PG_RGEN_INTERSECT_MNEE].cssRG);
unsigned int dss = std::max(stack_size[PG_CALL_SVM_AO].dssDC,
stack_size[PG_CALL_SVM_BEVEL].dssDC);
for (unsigned int i = 0; i < osl_stack_size.size(); ++i) {

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

@ -26,12 +26,12 @@ enum {
PG_RGEN_INTERSECT_SUBSURFACE,
PG_RGEN_INTERSECT_VOLUME_STACK,
PG_RGEN_INTERSECT_DEDICATED_LIGHT,
PG_RGEN_INTERSECT_MNEE,
PG_RGEN_SHADE_BACKGROUND,
PG_RGEN_SHADE_LIGHT_NEE,
PG_RGEN_SHADE_LIGHT_FORWARD,
PG_RGEN_SHADE_SURFACE,
PG_RGEN_SHADE_SURFACE_RAYTRACE,
PG_RGEN_SHADE_SURFACE_MNEE,
PG_RGEN_SHADE_VOLUME,
PG_RGEN_SHADE_VOLUME_RAY_MARCHING,
PG_RGEN_SHADE_SHADOW,

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

@ -121,9 +121,9 @@ bool OptiXDeviceQueue::enqueue(DeviceKernel kernel,
pipeline = optix_device->pipelines[PIP_SHADE];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SURFACE_RAYTRACE * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE:
pipeline = optix_device->pipelines[PIP_SHADE];
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_SHADE_SURFACE_MNEE * sizeof(SbtRecord);
sbt_params.raygenRecord = sbt_data_ptr + PG_RGEN_INTERSECT_MNEE * sizeof(SbtRecord);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME:
pipeline = optix_device->pipelines[PIP_SHADE];

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

@ -86,8 +86,6 @@ PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
integrator_shader_sort_counter_(device, "integrator_shader_sort_counter", MEM_READ_WRITE),
integrator_shader_raytrace_sort_counter_(
device, "integrator_shader_raytrace_sort_counter", MEM_READ_WRITE),
integrator_shader_mnee_sort_counter_(
device, "integrator_shader_mnee_sort_counter", MEM_READ_WRITE),
integrator_shader_sort_prefix_sum_(
device, "integrator_shader_sort_prefix_sum", MEM_READ_WRITE),
integrator_shader_sort_partition_key_offsets_(
@ -283,15 +281,6 @@ void PathTraceWorkGPU::alloc_integrator_sorting()
(int *)integrator_shader_raytrace_sort_counter_.device_pointer;
}
}
if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE) {
if (integrator_shader_mnee_sort_counter_.size() < sort_buckets) {
integrator_shader_mnee_sort_counter_.alloc(sort_buckets);
integrator_shader_mnee_sort_counter_.zero_to_device();
integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE] =
(int *)integrator_shader_mnee_sort_counter_.device_pointer;
}
}
}
}
@ -429,7 +418,13 @@ DeviceKernel PathTraceWorkGPU::get_most_queued_kernel() const
int max_num_queued = 0;
DeviceKernel kernel = DEVICE_KERNEL_NUM;
for (int i = 0; i < DEVICE_KERNEL_INTEGRATOR_NUM; i++) {
for (int i = 0; i < DEVICE_GPU_KERNEL_INTEGRATOR_NUM; i++) {
/* SHADOW_PATH_MNEE_PENDING is a sentinel marker on shadow slots holding an MNEE precompute
* payload; there is no kernel to dispatch for it. The slot transitions to a real shadow
* kernel (or terminates) when integrator_shade_surface runs on the main path. */
if (i == DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING) {
continue;
}
if (queue_counter->num_queued[i] > max_num_queued) {
kernel = (DeviceKernel)i;
max_num_queued = queue_counter->num_queued[i];
@ -453,11 +448,6 @@ void PathTraceWorkGPU::enqueue_reset()
{
queue_->zero_to_device(integrator_shader_raytrace_sort_counter_);
}
if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE &&
integrator_shader_mnee_sort_counter_.size() != 0)
{
queue_->zero_to_device(integrator_shader_mnee_sort_counter_);
}
/* Tiles enqueue need to know number of active paths, which is based on this counter. Zero the
* counter on the host side because `zero_to_device()` is not doing it. */
@ -472,7 +462,7 @@ bool PathTraceWorkGPU::enqueue_path_iteration()
const IntegratorQueueCounter *queue_counter = integrator_queue_counter_.data();
int num_active_paths = 0;
for (int i = 0; i < DEVICE_KERNEL_INTEGRATOR_NUM; i++) {
for (int i = 0; i < DEVICE_GPU_KERNEL_INTEGRATOR_NUM; i++) {
num_active_paths += queue_counter->num_queued[i];
}
@ -588,7 +578,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_DEDICATED_LIGHT: {
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE: {
/* Ray intersection kernels with integrator state. */
const DeviceKernelArguments args(&d_path_index, &work_size);
@ -601,7 +592,6 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
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_RAY_MARCHING:
case DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT: {
@ -731,7 +721,8 @@ void PathTraceWorkGPU::compact_shadow_paths()
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];
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW] +
queue_counter->num_queued[DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING];
/* Early out if there is nothing that needs to be compacted. */
if (num_active_paths == 0) {
@ -961,7 +952,7 @@ int PathTraceWorkGPU::num_active_main_paths_paths()
IntegratorQueueCounter *queue_counter = integrator_queue_counter_.data();
int num_paths = 0;
for (int i = 0; i < DEVICE_KERNEL_INTEGRATOR_NUM; i++) {
for (int i = 0; i < DEVICE_GPU_KERNEL_INTEGRATOR_NUM; i++) {
DCHECK_GE(queue_counter->num_queued[i], 0)
<< "Invalid number of queued states for kernel "
<< device_kernel_as_string(static_cast<DeviceKernel>(i));
@ -1317,15 +1308,14 @@ int PathTraceWorkGPU::shadow_catcher_count_possible_splits()
bool PathTraceWorkGPU::kernel_uses_sorting(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_SURFACE_RAYTRACE);
}
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_INTERSECT_MNEE ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME_RAY_MARCHING ||
kernel == DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT);
@ -1335,8 +1325,7 @@ bool PathTraceWorkGPU::kernel_creates_ao_paths(DeviceKernel kernel)
{
return (device_scene_->data.kernel_features & KERNEL_FEATURE_AO) &&
(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_SURFACE_RAYTRACE);
}
bool PathTraceWorkGPU::kernel_is_shadow_path(DeviceKernel kernel)

1
intern/cycles/integrator/path_trace_work_gpu.h
View file

@ -137,7 +137,6 @@ class PathTraceWorkGPU : public PathTraceWork {
/* Shader sorting. */
device_vector<int> integrator_shader_sort_counter_;
device_vector<int> integrator_shader_raytrace_sort_counter_;
device_vector<int> integrator_shader_mnee_sort_counter_;
device_vector<int> integrator_shader_sort_prefix_sum_;
device_vector<int> integrator_shader_sort_partition_key_offsets_;
/* Path split. */

1
intern/cycles/kernel/CMakeLists.txt
View file

@ -184,6 +184,7 @@ set(SRC_KERNEL_INTEGRATOR_HEADERS
integrator/init_from_camera.h
integrator/intersect_dedicated_light.h
integrator/intersect_closest.h
integrator/intersect_mnee.h
integrator/intersect_shadow.h
integrator/intersect_subsurface.h
integrator/intersect_volume_stack.h

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

@ -29,6 +29,7 @@
#include "kernel/integrator/init_from_camera.h"
#include "kernel/integrator/intersect_closest.h"
#include "kernel/integrator/intersect_dedicated_light.h"
#include "kernel/integrator/intersect_mnee.h"
#include "kernel/integrator/intersect_shadow.h"
#include "kernel/integrator/intersect_subsurface.h"
#include "kernel/integrator/intersect_volume_stack.h"
@ -227,6 +228,22 @@ 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_mnee,
const ccl_global int *path_index_array,
const int work_size)
{
# ifdef __MNEE__
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_mnee(nullptr, state));
}
# endif
}
ccl_gpu_kernel_postfix
# ifdef __KERNEL_ONEAPI__
# include "kernel/device/oneapi/context_intersect_end.h"
# endif
@ -344,21 +361,6 @@ 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_surface_mnee,
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_surface_mnee(nullptr, state, render_buffer));
}
}
ccl_gpu_kernel_postfix
# ifdef __KERNEL_ONEAPI__
# include "kernel/device/oneapi/context_intersect_end.h"
# endif

6
intern/cycles/kernel/device/hiprt/kernel.cpp
View file

@ -21,6 +21,7 @@
# include "kernel/integrator/intersect_closest.h"
# include "kernel/integrator/intersect_dedicated_light.h"
# include "kernel/integrator/intersect_mnee.h"
# include "kernel/integrator/intersect_shadow.h"
# include "kernel/integrator/intersect_subsurface.h"
# include "kernel/integrator/intersect_volume_stack.h"
@ -122,9 +123,8 @@ ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
}
ccl_gpu_kernel_postfix
ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
ccl_gpu_kernel_signature(integrator_shade_surface_mnee,
ccl_gpu_kernel_signature(integrator_intersect_mnee,
const ccl_global int *path_index_array,
ccl_global float *render_buffer,
const int work_size,
ccl_global hiprtGlobalStackBuffer stack_buffer)
{
@ -132,7 +132,7 @@ ccl_gpu_kernel_threads(GPU_HIPRT_KERNEL_BLOCK_NUM_THREADS)
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_shade_surface_mnee(kg, state, render_buffer));
ccl_gpu_kernel_call(integrator_intersect_mnee(kg, state));
}
}
ccl_gpu_kernel_postfix

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

@ -180,7 +180,7 @@ bool oneapi_kernel_is_required_for_features(const std::string &kernel_name,
}
if ((kernel_features & KERNEL_FEATURE_MNEE) == 0 &&
kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE)) !=
kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE)) !=
std::string::npos)
{
return false;
@ -200,6 +200,8 @@ bool oneapi_kernel_is_required_for_features(const std::string &kernel_name,
std::string::npos) ||
(kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE)) !=
std::string::npos) ||
(kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE)) !=
std::string::npos) ||
(kernel_name.find(device_kernel_as_string(
DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT)) != std::string::npos)))
{
@ -214,7 +216,7 @@ bool oneapi_kernel_is_compatible_with_hardware_raytracing(const std::string &ker
/* MNEE and Ray-trace kernels work correctly with Hardware Ray-tracing starting with Embree 4.1.
*/
# if defined(RTC_VERSION) && RTC_VERSION < 40100
return (kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE)) ==
return (kernel_name.find(device_kernel_as_string(DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE)) ==
std::string::npos) &&
(kernel_name.find(device_kernel_as_string(
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE)) == std::string::npos);
@ -427,6 +429,11 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
oneapi_kernel_integrator_intersect_dedicated_light);
break;
}
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE: {
oneapi_call(
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_intersect_mnee);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND: {
oneapi_call(
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_background);
@ -465,11 +472,6 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
oneapi_kernel_integrator_shade_surface_raytrace);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE: {
oneapi_call(
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_surface_mnee);
break;
}
case DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME: {
oneapi_call(
kg, cgh, global_size, local_size, args, oneapi_kernel_integrator_shade_volume);
@ -732,6 +734,7 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
/* Unsupported kernels */
case DEVICE_KERNEL_NUM:
case DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL:
case DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING:
kernel_assert(0);
break;
}

5
intern/cycles/kernel/device/optix/kernel_shader_raytrace.cu
View file

@ -7,6 +7,7 @@
#include "kernel/device/optix/kernel.cu"
#include "kernel/integrator/intersect_mnee.h"
#include "kernel/integrator/shade_surface.h"
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_raytrace()
@ -18,11 +19,11 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_raytr
integrator_shade_surface_raytrace(nullptr, path_index, kernel_params.render_buffer);
}
extern "C" __global__ void __raygen__kernel_optix_integrator_shade_surface_mnee()
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_mnee()
{
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_surface_mnee(nullptr, path_index, kernel_params.render_buffer);
integrator_intersect_mnee(nullptr, path_index);
}

3
intern/cycles/kernel/integrator/init_from_bake.h
View file

@ -324,8 +324,7 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
const bool use_raytrace_kernel = (shader_flags & SD_HAS_RAYTRACE);
if (use_caustics) {
integrator_path_init_sorted(
kg, state, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE, shader_index);
integrator_path_init(state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
}
else if (use_raytrace_kernel) {
integrator_path_init_sorted(

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

@ -147,11 +147,11 @@ ccl_device_forceinline void integrator_split_shadow_catcher(
const int shader = intersection_get_shader(kg, isect);
const int flags = kernel_data_fetch(shaders, shader).flags;
const bool use_caustics = kernel_data.integrator.use_caustics &&
(object_flags & SD_OBJECT_CAUSTICS);
(object_flags & SD_OBJECT_CAUSTICS_RECEIVER);
const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
if (use_caustics) {
integrator_path_init_sorted(kg, state, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE, shader);
integrator_path_init(state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
}
else if (use_raytrace_kernel) {
integrator_path_init_sorted(
@ -176,12 +176,11 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_shadow_catche
const int flags = kernel_data_fetch(shaders, shader).flags;
const uint object_flags = intersection_get_object_flags(kg, &isect);
const bool use_caustics = kernel_data.integrator.use_caustics &&
(object_flags & SD_OBJECT_CAUSTICS);
(object_flags & SD_OBJECT_CAUSTICS_RECEIVER);
const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
if (use_caustics) {
integrator_path_next_sorted(
kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE, shader);
integrator_path_next(state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
}
else if (use_raytrace_kernel) {
integrator_path_next_sorted(
@ -261,11 +260,10 @@ ccl_device_forceinline void integrator_intersect_next_kernel(
if (!integrator_intersect_terminate(kg, state, flags)) {
const uint object_flags = intersection_get_object_flags(kg, isect);
const bool use_caustics = kernel_data.integrator.use_caustics &&
(object_flags & SD_OBJECT_CAUSTICS);
(object_flags & SD_OBJECT_CAUSTICS_RECEIVER);
const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
if (use_caustics) {
integrator_path_next_sorted(
kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE, shader);
integrator_path_next(state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
}
else if (use_raytrace_kernel) {
integrator_path_next_sorted(
@ -320,12 +318,11 @@ ccl_device_forceinline void integrator_intersect_next_kernel_after_volume(
const int flags = kernel_data_fetch(shaders, shader).flags;
const uint object_flags = intersection_get_object_flags(kg, isect);
const bool use_caustics = kernel_data.integrator.use_caustics &&
(object_flags & SD_OBJECT_CAUSTICS);
(object_flags & SD_OBJECT_CAUSTICS_RECEIVER);
const bool use_raytrace_kernel = (flags & SD_HAS_RAYTRACE);
if (use_caustics) {
integrator_path_next_sorted(
kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE, shader);
integrator_path_next(state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
}
else if (use_raytrace_kernel) {
integrator_path_next_sorted(

119
intern/cycles/kernel/integrator/intersect_mnee.h Normal file
View file

@ -0,0 +1,119 @@
/* SPDX-FileCopyrightText: 2011-2026 Blender Foundation
*
* SPDX-License-Identifier: Apache-2.0 */
#pragma once
#include "kernel/integrator/mnee.h"
#include "kernel/integrator/shade_surface.h"
CCL_NAMESPACE_BEGIN
#ifdef __MNEE__
/* Sample a light and run the MNEE manifold walk for a caustic receiver. On a successful walk
* the result is written to a shadow slot for integrator_shade_surface, otherwise nothing is
* written and direct light is sampled there. */
ccl_device_forceinline ShaderEvalResult
integrate_surface_mnee(KernelGlobals kg,
IntegratorState state,
ccl_private ShaderData *sd,
const ccl_private RNGState *rng_state)
{
/* Kernel must only be scheduled for caustic receivers. */
kernel_assert(sd->object_flag & SD_OBJECT_CAUSTICS_RECEIVER);
if (!kernel_data.integrator.use_direct_light) {
return SHADER_EVAL_OK;
}
/* Sample position on a light. */
LightSample ls ccl_optional_struct_init;
{
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
const uint bounce = INTEGRATOR_STATE(state, path, bounce);
const float3 rand_light = path_state_rng_3D(kg, rng_state, PRNG_LIGHT);
if (!light_sample_from_position(kg,
rand_light,
sd->time,
sd->P,
sd->N,
light_link_receiver_nee(kg, sd),
sd->flag,
bounce,
path_flag,
&ls))
{
return SHADER_EVAL_OK;
}
}
kernel_assert(ls.pdf != 0.0f);
/* The manifold walk connects a caustic light to the receiver across reflection; transmission
* caustics and triangle lights are not handled. */
if (ls.type == LIGHT_TRIANGLE || dot(ls.D, sd->N) < 0.0f) {
return SHADER_EVAL_OK;
}
if (!kernel_data_fetch(lights, ls.prim).use_caustics) {
return SHADER_EVAL_OK;
}
ShaderDataCausticsStorage emission_sd_storage;
ccl_private ShaderData *emission_sd = AS_SHADER_DATA(&emission_sd_storage);
Spectrum mnee_throughput = zero_spectrum();
float3 mnee_wo = zero_float3();
int mnee_vertex_count = 0;
const ShaderEvalResult result = kernel_path_mnee_sample(
kg, state, sd, emission_sd, rng_state, &ls, &mnee_throughput, &mnee_wo, mnee_vertex_count);
if (result == SHADER_EVAL_CACHE_MISS) {
return SHADER_EVAL_CACHE_MISS;
}
/* Store MNEE state in a shadow state, to avoid increasing path state size.
* This is then turned into an actual shadow ray state in shade_surface, or discarded. */
if (mnee_vertex_count > 0) {
Ray ray ccl_optional_struct_init;
light_sample_to_surface_shadow_ray(kg, emission_sd, &ls, &ray);
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg, state, DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING, false);
integrator_state_write_mnee(
state, shadow_state, &ls, &ray, mnee_vertex_count, mnee_throughput, mnee_wo);
}
return SHADER_EVAL_OK;
}
#endif /* __MNEE__ */
ccl_device void integrator_intersect_mnee(KernelGlobals kg, IntegratorState state)
{
PROFILING_INIT(kg, PROFILING_SHADE_SURFACE_DIRECT_LIGHT);
ShaderData sd;
integrate_surface_shader_setup(kg, state, &sd);
const int shader = sd.shader & SHADER_MASK;
#ifdef __MNEE__
RNGState rng_state;
path_state_rng_load(state, &rng_state);
const ShaderEvalResult result = integrate_surface_mnee(kg, state, &sd, &rng_state);
if (result == SHADER_EVAL_CACHE_MISS) {
integrator_path_cache_miss(state, DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
return;
}
#endif
integrator_path_next_sorted(kg,
state,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE,
shader);
}
CCL_NAMESPACE_END

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

@ -6,6 +6,7 @@
#include "kernel/integrator/intersect_closest.h"
#include "kernel/integrator/intersect_dedicated_light.h"
#include "kernel/integrator/intersect_mnee.h"
#include "kernel/integrator/intersect_shadow.h"
#include "kernel/integrator/intersect_subsurface.h"
#include "kernel/integrator/intersect_volume_stack.h"
@ -32,18 +33,21 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
switch (shadow_queued_kernel) {
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW:
integrator_intersect_shadow(kg, &state->shadow);
break;
continue;
case DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW:
integrator_shade_shadow(kg, &state->shadow, render_buffer);
break;
continue;
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE:
integrator_shade_light_nee(kg, &state->shadow, render_buffer);
continue;
case DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING:
/* Not a real kernel, only a state to keep it alive until
* shade_surface uses this shadow path. */
break;
default:
kernel_assert(0);
break;
}
continue;
}
/* Handle any AO paths before we potentially create more AO paths. */
@ -85,9 +89,6 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE:
integrator_shade_surface_raytrace(kg, state, render_buffer);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE:
integrator_shade_surface_mnee(kg, state, render_buffer);
break;
case DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_FORWARD:
integrator_shade_light_forward(kg, state, render_buffer);
break;
@ -103,6 +104,9 @@ ccl_device void integrator_megakernel(KernelGlobals kg,
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT:
integrator_intersect_dedicated_light(kg, state);
break;
case DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE:
integrator_intersect_mnee(kg, state);
break;
default:
kernel_assert(0);
break;

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

@ -796,13 +796,15 @@ ccl_device_inline ShaderEvalResult mnee_path_contribution(KernelGlobals kg,
const bool light_fixed_direction,
const int vertex_count,
ccl_private ManifoldVertex *vertices,
ccl_private BsdfEval *throughput)
ccl_private Spectrum *throughput,
ccl_private float3 *r_receiver_wo)
{
float wo_len;
float3 wo = normalize_len(vertices[0].p - sd->P, &wo_len);
/* Initialize throughput and evaluate receiver bsdf * |n.wo|. */
surface_shader_bsdf_eval(kg, state, sd, wo, throughput, ls->shader);
/* Initialize throughput. */
*r_receiver_wo = wo;
*throughput = one_spectrum();
/* Update light sample with new position / direction and keep pdf in vertex area measure. */
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
@ -831,7 +833,7 @@ ccl_device_inline ShaderEvalResult mnee_path_contribution(KernelGlobals kg,
return SHADER_EVAL_CACHE_MISS;
}
bsdf_eval_mul(throughput, ls->eval_fac / ls->pdf);
*throughput *= ls->eval_fac / ls->pdf;
/* Generalized geometry term. */
float dh_dx;
@ -842,12 +844,12 @@ ccl_device_inline ShaderEvalResult mnee_path_contribution(KernelGlobals kg,
return SHADER_EVAL_EMPTY;
}
/* Receiver bsdf eval above already contains |n.wo|. */
/* Receiver bsdf eval in shade_surface already contains |n.wo|. */
const float dw0_dx1 = fabsf(dot(wo, vertices[0].n)) / sqr(wo_len);
/* Clamp since it has a tendency to be unstable. */
const float G = fminf(dw0_dx1 * dx1_dxlight, 2.f);
bsdf_eval_mul(throughput, G);
*throughput *= G;
/* Specular reflectance. */
@ -930,7 +932,7 @@ ccl_device_inline ShaderEvalResult mnee_path_contribution(KernelGlobals kg,
* divided by corresponding sampled pdf:
* fr(vi)_do / pdf_dh(vi) x |do/dh| x |n.wo / n.h| */
const Spectrum bsdf_contribution = mnee_eval_bsdf_contribution(kg, v.bsdf, wi, wo);
bsdf_eval_mul(throughput, bsdf_contribution);
*throughput *= bsdf_contribution;
}
/* Restore original state path bounce info. */
@ -948,7 +950,8 @@ ccl_device_inline ShaderEvalResult kernel_path_mnee_sample(KernelGlobals kg,
ccl_private ShaderData *sd_mnee,
const ccl_private RNGState *rng_state,
ccl_private LightSample *ls,
ccl_private BsdfEval *throughput,
ccl_private Spectrum *throughput,
ccl_private float3 *r_receiver_wo,
ccl_private int &r_vertex_count)
{
/*
@ -1109,8 +1112,16 @@ ccl_device_inline ShaderEvalResult kernel_path_mnee_sample(KernelGlobals kg,
* each interface. */
if (mnee_newton_solver(kg, sd, sd_mnee, ls, light_fixed_direction, vertex_count, vertices)) {
/* 3. If a solution exists, calculate contribution of the corresponding path */
ShaderEvalResult result = mnee_path_contribution(
kg, state, sd, sd_mnee, ls, light_fixed_direction, vertex_count, vertices, throughput);
ShaderEvalResult result = mnee_path_contribution(kg,
state,
sd,
sd_mnee,
ls,
light_fixed_direction,
vertex_count,
vertices,
throughput,
r_receiver_wo);
/* TODO: Cache misses are not handled correctly.
* - PATH_MNEE_VALID flag is not handled properly
* - AOVs and other passes have already been written at this point

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

@ -16,8 +16,6 @@
#include "kernel/geom/motion_triangle.h"
#include "kernel/geom/triangle.h"
#include "kernel/integrator/mnee.h"
#include "kernel/integrator/guiding.h"
#include "kernel/integrator/shadow_linking.h"
#include "kernel/integrator/subsurface.h"
@ -213,14 +211,24 @@ integrate_direct_light_shadow_init_common(KernelGlobals kg,
const int mnee_vertex_count,
const bool constant_light_shader)
{
const DeviceKernel next_kernel = (constant_light_shader) ?
DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW :
DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE;
/* Branch off shadow kernel. */
IntegratorShadowState shadow_state = integrator_shadow_path_init(
kg,
state,
(constant_light_shader) ? DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW :
DEVICE_KERNEL_INTEGRATOR_SHADE_LIGHT_NEE,
false);
IntegratorShadowState shadow_state;
#ifdef __MNEE__
if (mnee_vertex_count > 0) {
/* Reuse shadow path that was already allocated by intersect_mnee. */
shadow_state = integrator_state_get_mnee_shadow_state(state);
integrator_shadow_path_next(
shadow_state, DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING, next_kernel);
}
else
#endif
{
shadow_state = integrator_shadow_path_init(kg, state, next_kernel, false);
}
#ifdef __VOLUME__
/* Copy volume stack and enter/exit volume. */
@ -318,9 +326,20 @@ ccl_device
return SHADER_EVAL_EMPTY;
}
/* Sample position on a light. */
LightSample ls ccl_optional_struct_init;
int mnee_vertex_count = 0; // NOLINT
#ifdef __MNEE__
if ((kernel_data.kernel_features & KERNEL_FEATURE_MNEE) &&
(INTEGRATOR_STATE(state, path, mnee) & PATH_MNEE_SAMPLED))
{
/* MNEE already sampled a light and caustics casters. */
integrator_state_read_mnee(state, &ls, &mnee_vertex_count);
}
else
#endif
{
/* Sample position on a light. */
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
const uint bounce = INTEGRATOR_STATE(state, path, bounce);
const float3 rand_light = path_state_rng_3D(kg, rng_state, PRNG_LIGHT);
@ -352,41 +371,15 @@ ccl_device
}
}
Ray ray ccl_optional_struct_init;
BsdfEval bsdf_eval ccl_optional_struct_init;
int mnee_vertex_count = 0; // NOLINT
#ifdef __MNEE__
IF_KERNEL_FEATURE(MNEE)
{
if (ls.type != LIGHT_TRIANGLE) {
/* Is this a caustic light? */
const bool use_caustics = kernel_data_fetch(lights, ls.prim).use_caustics;
if (use_caustics) {
/* Are we on a caustic caster? */
if (is_transmission && (sd->object_flag & SD_OBJECT_CAUSTICS_CASTER)) {
return SHADER_EVAL_EMPTY;
}
/* 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);
ShaderEvalResult result = kernel_path_mnee_sample(
kg, state, sd, emission_sd, rng_state, &ls, &bsdf_eval, mnee_vertex_count);
if (result == SHADER_EVAL_CACHE_MISS) {
return SHADER_EVAL_CACHE_MISS;
}
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);
}
}
}
/* On a caustic caster, a caustic light's contribution is delivered to receivers by
* MNEE and does not need to be computed again here. */
if (kernel_data.kernel_features & KERNEL_FEATURE_MNEE) {
if (mnee_vertex_count == 0 && is_transmission &&
(sd->object_flag & SD_OBJECT_CAUSTICS_CASTER) && ls.type != LIGHT_TRIANGLE &&
kernel_data_fetch(lights, ls.prim).use_caustics)
{
return SHADER_EVAL_EMPTY;
}
}
#endif
@ -396,15 +389,26 @@ ccl_device
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 bsdf_eval ccl_optional_struct_init;
const float bsdf_pdf = surface_shader_bsdf_eval(kg, state, sd, ls.D, &bsdf_eval, ls.shader);
Ray ray ccl_optional_struct_init;
#ifdef __MNEE__
if (mnee_vertex_count > 0) {
light_shader_eval *= integrator_state_read_mnee_throughput(state);
bsdf_eval_mul(&bsdf_eval, light_shader_eval);
if (bsdf_eval_is_zero(&bsdf_eval)) {
return SHADER_EVAL_EMPTY;
}
integrator_state_read_mnee_ray(state, &ls, &ray);
}
else
#endif /* __MNEE__ */
{
/* 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_shader_eval * ls.eval_fac / ls.pdf * mis_weight);
@ -894,6 +898,24 @@ ccl_device_forceinline void integrator_shade_surface(KernelGlobals kg,
integrator_path_cache_miss_sorted(state, current_kernel);
return;
}
#ifdef __MNEE__
/* Cleanup MNEE flag and shadow path if it was not reused for shadow trace. */
if ((kernel_data.kernel_features & KERNEL_FEATURE_MNEE) &&
(INTEGRATOR_STATE(state, path, mnee) & PATH_MNEE_SAMPLED))
{
INTEGRATOR_STATE_WRITE(state, path, mnee) &= ~PATH_MNEE_SAMPLED;
const IntegratorShadowState shadow_state = integrator_state_get_mnee_shadow_state(state);
if (INTEGRATOR_STATE(shadow_state, shadow_path, queued_kernel) ==
DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING)
{
integrator_shadow_path_terminate(shadow_state,
DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING);
}
}
#endif
if (continue_path_label == LABEL_NONE) {
integrator_path_terminate(kg, state, render_buffer, current_kernel);
return;
@ -920,14 +942,4 @@ ccl_device_forceinline void integrator_shade_surface_raytrace(
kg, state, render_buffer);
}
ccl_device_forceinline void integrator_shade_surface_mnee(
KernelGlobals kg, IntegratorState state, ccl_global float *ccl_restrict render_buffer)
{
#ifdef __MNEE__
integrator_shade_surface<(KERNEL_FEATURE_NODE_MASK_SURFACE & ~KERNEL_FEATURE_NODE_RAYTRACE) |
KERNEL_FEATURE_MNEE,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE>(kg, state, render_buffer);
#endif
}
CCL_NAMESPACE_END

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

@ -98,7 +98,7 @@ struct IntegratorStateCPU {
* Keep track of which kernels are queued to be executed next in the path
* for GPU rendering. */
struct IntegratorQueueCounter {
int num_queued[DEVICE_KERNEL_INTEGRATOR_NUM];
int num_queued[DEVICE_GPU_KERNEL_INTEGRATOR_NUM];
int cache_miss;
};
@ -197,7 +197,7 @@ struct IntegratorStateGPU {
ccl_global IntegratorQueueCounter *queue_counter;
/* Count number of kernels queued for specific shaders. */
ccl_global int *sort_key_counter[DEVICE_KERNEL_INTEGRATOR_NUM];
ccl_global int *sort_key_counter[DEVICE_GPU_KERNEL_INTEGRATOR_NUM];
/* Index of shadow path which will be used by a next shadow path. */
ccl_global int *next_shadow_path_index;

2
intern/cycles/kernel/integrator/state_template.h
View file

@ -41,6 +41,8 @@ KERNEL_STRUCT_MEMBER(path, uint8_t, visibility, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(path, uint32_t, flag, KERNEL_FEATURE_PATH_TRACING)
/* enum PathRayMNEE */
KERNEL_STRUCT_MEMBER(path, uint8_t, mnee, KERNEL_FEATURE_PATH_TRACING)
/* Index of shadow state path used for storing MNEE state. */
KERNEL_STRUCT_MEMBER(path, int, mnee_shadow_state, KERNEL_FEATURE_MNEE)
/* Majorant volume optical depth. */
KERNEL_STRUCT_MEMBER(path, float, optical_depth, KERNEL_FEATURE_PATH_TRACING)
/* Multiple importance sampling

161
intern/cycles/kernel/integrator/state_util.h
View file

@ -8,6 +8,8 @@
#include "kernel/integrator/state.h"
#include "kernel/light/common.h"
#include "kernel/sample/lcg.h"
#include "kernel/util/differential.h"
@ -309,6 +311,149 @@ ccl_device_forceinline void integrator_state_read_shadow_isect(
isect->t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, index, t);
}
/* MNEE state.
*
* This is packed into the shadow_state to avoid increasing overall path state size. */
#ifdef __MNEE__
ccl_device_forceinline IntegratorShadowState
integrator_state_get_mnee_shadow_state(ConstIntegratorState state)
{
# ifdef __KERNEL_GPU__
return IntegratorShadowState(INTEGRATOR_STATE(state, path, mnee_shadow_state));
# else
return &(((IntegratorStateCPU *)state)->shadow);
# endif
}
# ifdef __KERNEL_GPU__
/* The MNEE shadow slot stores a reference to its owning main path so shadow path
* sorting can maintain the correct index. */
ccl_device_forceinline void integrator_state_write_mnee_shadow_owner(
IntegratorShadowState shadow_state, IntegratorState state)
{
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 2, prim) = (int)state;
}
ccl_device_forceinline IntegratorState
integrator_state_read_mnee_shadow_owner(ConstIntegratorShadowState shadow_state)
{
return IntegratorState(INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 2, prim));
}
# endif
ccl_device_forceinline void integrator_state_write_mnee(IntegratorState state,
IntegratorShadowState shadow_state,
const ccl_private LightSample *ls,
const ccl_private Ray *ray,
const int mnee_vertex_count,
const Spectrum mnee_throughput,
const float3 mnee_wo)
{
static_assert(INTEGRATOR_SHADOW_ISECT_SIZE >= 2);
# ifdef __KERNEL_GPU__
INTEGRATOR_STATE_WRITE(state, path, mnee_shadow_state) = (int)shadow_state;
integrator_state_write_mnee_shadow_owner(shadow_state, state);
# endif
/* Light sample. */
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, t) = ls->P.x;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, u) = ls->P.y;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, v) = ls->P.z;
/* When the integrate_surface_direct_light() reads the MNEE state it should read mnee_wo as the
* light sample direction. */
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, t) = mnee_wo.x;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, u) = mnee_wo.y;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, v) = mnee_wo.z;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, tmin) = ls->t;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, tmax) = ls->pdf;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, time) = ls->eval_fac;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, self_light_object) = ls->object;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, self_light_prim) = ls->prim;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, object) = ls->shader;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, prim) = ls->group + 1;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 0, type) = (int)ls->type;
/* Ray. */
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, prim) = mnee_vertex_count;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, throughput) = mnee_throughput;
/* The ray direction becomes the original light sample's direction for the shadow ray tracing. */
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, D) = ls->D;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, P) = ray->P;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_ray, dP) = ray->dP;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, object) = ray->self.object;
INTEGRATOR_STATE_ARRAY_WRITE(shadow_state, shadow_isect, 1, type) = ray->self.prim;
INTEGRATOR_STATE_WRITE(state, path, mnee) |= PATH_MNEE_SAMPLED;
}
ccl_device_forceinline void integrator_state_read_mnee(ConstIntegratorState state,
ccl_private LightSample *ls,
ccl_private int *mnee_vertex_count)
{
static_assert(INTEGRATOR_SHADOW_ISECT_SIZE >= 2);
ConstIntegratorShadowState shadow_state = integrator_state_get_mnee_shadow_state(state);
ls->P = make_float3(INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 0, t),
INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 0, u),
INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 0, v));
ls->D = make_float3(INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 1, t),
INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 1, u),
INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 1, v));
ls->t = INTEGRATOR_STATE(shadow_state, shadow_ray, tmin);
ls->pdf = INTEGRATOR_STATE(shadow_state, shadow_ray, tmax);
ls->eval_fac = INTEGRATOR_STATE(shadow_state, shadow_ray, time);
ls->object = INTEGRATOR_STATE(shadow_state, shadow_ray, self_light_object);
ls->prim = INTEGRATOR_STATE(shadow_state, shadow_ray, self_light_prim);
ls->shader = INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 0, object);
ls->group = INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 0, prim) - 1;
ls->type = (LightType)INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 0, type);
ls->pdf_selection = 0.0f;
ls->emitter_id = EMITTER_NONE;
*mnee_vertex_count = INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 1, prim);
}
ccl_device_forceinline void integrator_state_read_mnee_ray(ConstIntegratorState state,
const ccl_private LightSample *ls,
ccl_private Ray *ray)
{
static_assert(INTEGRATOR_SHADOW_ISECT_SIZE >= 2);
ConstIntegratorShadowState shadow_state = integrator_state_get_mnee_shadow_state(state);
ray->P = INTEGRATOR_STATE(shadow_state, shadow_ray, P);
if (ls->t == FLT_MAX) {
/* Distant light. */
ray->D = INTEGRATOR_STATE(shadow_state, shadow_ray, D);
ray->tmax = ls->t;
}
else {
/* Other lights. */
ray->D = ls->P - ray->P;
ray->D = safe_normalize_len(ray->D, &ray->tmax);
}
ray->tmin = ((ls->shader & SHADER_CAST_SHADOW) == 0) ? FLT_MAX : 0.0f;
ray->time = INTEGRATOR_STATE(state, ray, time);
ray->dP = INTEGRATOR_STATE(shadow_state, shadow_ray, dP);
ray->dD = differential_zero_compact();
ray->self.object = INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 1, object);
ray->self.prim = INTEGRATOR_STATE_ARRAY(shadow_state, shadow_isect, 1, type);
ray->self.light_object = ls->object;
ray->self.light_prim = ls->prim;
}
ccl_device_forceinline Spectrum integrator_state_read_mnee_throughput(ConstIntegratorState state)
{
ConstIntegratorShadowState shadow_state = integrator_state_get_mnee_shadow_state(state);
return INTEGRATOR_STATE(shadow_state, shadow_path, throughput);
}
#endif /* __MNEE__ */
#if defined(__KERNEL_GPU__)
ccl_device_inline void integrator_state_copy_only(KernelGlobals kg,
ConstIntegratorState to_state,
@ -376,6 +521,13 @@ ccl_device_inline void integrator_state_move(KernelGlobals kg,
integrator_state_copy_only(kg, to_state, state);
INTEGRATOR_STATE_WRITE(state, path, queued_kernel) = 0;
# ifdef __MNEE__
if (INTEGRATOR_STATE(to_state, path, mnee) & PATH_MNEE_SAMPLED) {
const IntegratorShadowState slot = INTEGRATOR_STATE(to_state, path, mnee_shadow_state);
integrator_state_write_mnee_shadow_owner(slot, to_state);
}
# endif
}
ccl_device_inline void integrator_shadow_state_copy_only(KernelGlobals kg,
@ -444,6 +596,15 @@ ccl_device_inline void integrator_shadow_state_move(KernelGlobals kg,
integrator_shadow_state_copy_only(kg, to_state, state);
INTEGRATOR_STATE_WRITE(state, shadow_path, queued_kernel) = 0;
# ifdef __MNEE__
if (INTEGRATOR_STATE(to_state, shadow_path, queued_kernel) ==
DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING)
{
const IntegratorState main_state = integrator_state_read_mnee_shadow_owner(to_state);
INTEGRATOR_STATE_WRITE(main_state, path, mnee_shadow_state) = (int)to_state;
}
# endif
}
#endif

8
intern/cycles/kernel/integrator/subsurface.h
View file

@ -222,11 +222,9 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat
const bool use_raytrace_kernel = (shader_flags & SD_HAS_RAYTRACE);
if (use_caustics) {
integrator_path_next_sorted(kg,
state,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE,
DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE,
shader);
integrator_path_next(state,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE);
}
else if (use_raytrace_kernel) {
integrator_path_next_sorted(kg,

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

@ -280,6 +280,9 @@ enum PathRayMNEE {
PATH_MNEE_VALID = (1U << 0U),
PATH_MNEE_RECEIVER_ANCESTOR = (1U << 1U),
PATH_MNEE_CULL_LIGHT_CONNECTION = (1U << 2U),
/* MNEE path was successfully sampled in intersect_mnee. */
PATH_MNEE_SAMPLED = (1U << 3U),
};
/* Configure ray visibility bits for rays and objects respectively,
@ -1746,16 +1749,17 @@ enum DeviceKernel : int {
DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT,
DEVICE_KERNEL_INTEGRATOR_INTERSECT_MNEE,
DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND,
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,
DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME,
DEVICE_KERNEL_INTEGRATOR_SHADE_VOLUME_RAY_MARCHING,
DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW,
DEVICE_KERNEL_INTEGRATOR_SHADE_DEDICATED_LIGHT,
DEVICE_KERNEL_INTEGRATOR_SHADOW_PATH_MNEE_PENDING,
DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL,
DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY,
@ -1819,7 +1823,8 @@ enum DeviceKernel : int {
};
enum {
DEVICE_KERNEL_INTEGRATOR_NUM = DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL + 1,
/* Megakernel is the first kernel not used by GPU integrator. */
DEVICE_GPU_KERNEL_INTEGRATOR_NUM = DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL,
};
CCL_NAMESPACE_END

8
tests/python/cycles_render_tests.py
View file

@ -108,11 +108,6 @@ if platform.system() == "Darwin":
"underwater_caustics.blend",
]
BLOCKLIST_HIP = [
# MNEE does not work properly on RDNA2 GPUs.
'underwater_caustics.blend',
]
BLOCKLIST_GPU = [
# Uninvestigated differences with GPU.
'glass_mix_40964.blend',
@ -304,9 +299,6 @@ def main():
blocklist += BLOCKLIST_METAL
blocklist += BLOCKLIST_METAL_RT
if device in ('HIP', 'HIP-RT'):
blocklist += BLOCKLIST_HIP
test_dir_name = Path(args.testdir).name
report = CyclesReport('Cycles', test_dir_name, args.outdir, args.oiiotool, device, blocklist, args.osl == 'all')