diff options
Diffstat (limited to 'intern/cycles/kernel/kernels')
9 files changed, 186 insertions, 0 deletions
diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h index f5d981fb71a..683f4b88d79 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu.h @@ -89,5 +89,9 @@ DECLARE_SPLIT_KERNEL_FUNCTION(enqueue_inactive) DECLARE_SPLIT_KERNEL_FUNCTION(next_iteration_setup) DECLARE_SPLIT_KERNEL_FUNCTION(indirect_subsurface) DECLARE_SPLIT_KERNEL_FUNCTION(buffer_update) +DECLARE_SPLIT_KERNEL_FUNCTION(adaptive_stopping) +DECLARE_SPLIT_KERNEL_FUNCTION(adaptive_filter_x) +DECLARE_SPLIT_KERNEL_FUNCTION(adaptive_filter_y) +DECLARE_SPLIT_KERNEL_FUNCTION(adaptive_adjust_samples) #undef KERNEL_ARCH diff --git a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h index 9ca3f46b5b6..96b2bf11132 100644 --- a/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h +++ b/intern/cycles/kernel/kernels/cpu/kernel_cpu_impl.h @@ -58,6 +58,10 @@ # include "kernel/split/kernel_next_iteration_setup.h" # include "kernel/split/kernel_indirect_subsurface.h" # include "kernel/split/kernel_buffer_update.h" +# include "kernel/split/kernel_adaptive_stopping.h" +# include "kernel/split/kernel_adaptive_filter_x.h" +# include "kernel/split/kernel_adaptive_filter_y.h" +# include "kernel/split/kernel_adaptive_adjust_samples.h" # endif /* __SPLIT_KERNEL__ */ #else # define STUB_ASSERT(arch, name) \ @@ -204,6 +208,10 @@ DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(enqueue_inactive, uint) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(next_iteration_setup, uint) DEFINE_SPLIT_KERNEL_FUNCTION(indirect_subsurface) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(buffer_update, uint) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_stopping) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_filter_x) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_filter_y) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_adjust_samples) #endif /* __SPLIT_KERNEL__ */ #undef KERNEL_STUB diff --git a/intern/cycles/kernel/kernels/cuda/kernel.cu b/intern/cycles/kernel/kernels/cuda/kernel.cu index af311027f78..c4c810c6a82 100644 --- a/intern/cycles/kernel/kernels/cuda/kernel.cu +++ b/intern/cycles/kernel/kernels/cuda/kernel.cu @@ -33,6 +33,7 @@ #include "kernel/kernel_path_branched.h" #include "kernel/kernel_bake.h" #include "kernel/kernel_work_stealing.h" +#include "kernel/kernel_adaptive_sampling.h" /* kernels */ extern "C" __global__ void @@ -83,6 +84,75 @@ kernel_cuda_branched_path_trace(WorkTile *tile, uint total_work_size) extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) +kernel_cuda_adaptive_stopping(WorkTile *tile, int sample, uint total_work_size) +{ + int work_index = ccl_global_id(0); + bool thread_is_active = work_index < total_work_size; + KernelGlobals kg; + if(thread_is_active && kernel_data.film.pass_adaptive_aux_buffer) { + uint x = tile->x + work_index % tile->w; + uint y = tile->y + work_index / tile->w; + int index = tile->offset + x + y * tile->stride; + ccl_global float *buffer = tile->buffer + index * kernel_data.film.pass_stride; + kernel_do_adaptive_stopping(&kg, buffer, sample); + } +} + +extern "C" __global__ void +CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) +kernel_cuda_adaptive_filter_x(WorkTile *tile, int sample, uint) +{ + KernelGlobals kg; + if(kernel_data.film.pass_adaptive_aux_buffer && sample > kernel_data.integrator.adaptive_min_samples) { + if(ccl_global_id(0) < tile->h) { + int y = tile->y + ccl_global_id(0); + kernel_do_adaptive_filter_x(&kg, y, tile); + } + } +} + +extern "C" __global__ void +CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) +kernel_cuda_adaptive_filter_y(WorkTile *tile, int sample, uint) +{ + KernelGlobals kg; + if(kernel_data.film.pass_adaptive_aux_buffer && sample > kernel_data.integrator.adaptive_min_samples) { + if(ccl_global_id(0) < tile->w) { + int x = tile->x + ccl_global_id(0); + kernel_do_adaptive_filter_y(&kg, x, tile); + } + } +} + +extern "C" __global__ void +CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) +kernel_cuda_adaptive_scale_samples(WorkTile *tile, int start_sample, int sample, uint total_work_size) +{ + if(kernel_data.film.pass_adaptive_aux_buffer) { + int work_index = ccl_global_id(0); + bool thread_is_active = work_index < total_work_size; + KernelGlobals kg; + if(thread_is_active) { + uint x = tile->x + work_index % tile->w; + uint y = tile->y + work_index / tile->w; + int index = tile->offset + x + y * tile->stride; + ccl_global float *buffer = tile->buffer + index * kernel_data.film.pass_stride; + if(buffer[kernel_data.film.pass_sample_count] < 0.0f) { + buffer[kernel_data.film.pass_sample_count] = -buffer[kernel_data.film.pass_sample_count]; + float sample_multiplier = sample / max((float)start_sample + 1.0f, buffer[kernel_data.film.pass_sample_count]); + if(sample_multiplier != 1.0f) { + kernel_adaptive_post_adjust(&kg, buffer, sample_multiplier); + } + } + else { + kernel_adaptive_post_adjust(&kg, buffer, sample / (sample - 1.0f)); + } + } + } +} + +extern "C" __global__ void +CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) kernel_cuda_convert_to_byte(uchar4 *rgba, float *buffer, float sample_scale, int sx, int sy, int sw, int sh, int offset, int stride) { int x = sx + blockDim.x*blockIdx.x + threadIdx.x; diff --git a/intern/cycles/kernel/kernels/cuda/kernel_split.cu b/intern/cycles/kernel/kernels/cuda/kernel_split.cu index 43b3d0aa0e6..95ad7599cf1 100644 --- a/intern/cycles/kernel/kernels/cuda/kernel_split.cu +++ b/intern/cycles/kernel/kernels/cuda/kernel_split.cu @@ -43,6 +43,10 @@ #include "kernel/split/kernel_next_iteration_setup.h" #include "kernel/split/kernel_indirect_subsurface.h" #include "kernel/split/kernel_buffer_update.h" +#include "kernel/split/kernel_adaptive_stopping.h" +#include "kernel/split/kernel_adaptive_filter_x.h" +#include "kernel/split/kernel_adaptive_filter_y.h" +#include "kernel/split/kernel_adaptive_adjust_samples.h" #include "kernel/kernel_film.h" @@ -121,6 +125,10 @@ DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(enqueue_inactive, uint) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(next_iteration_setup, uint) DEFINE_SPLIT_KERNEL_FUNCTION(indirect_subsurface) DEFINE_SPLIT_KERNEL_FUNCTION_LOCALS(buffer_update, uint) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_stopping) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_filter_x) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_filter_y) +DEFINE_SPLIT_KERNEL_FUNCTION(adaptive_adjust_samples) extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS) diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl new file mode 100644 index 00000000000..ebdb99d4730 --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl @@ -0,0 +1,23 @@ +/* + * Copyright 2019 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "kernel/kernel_compat_opencl.h" +#include "kernel/split/kernel_split_common.h" +#include "kernel/split/kernel_adaptive_adjust_samples.h" + +#define KERNEL_NAME adaptive_adjust_samples +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_x.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_x.cl new file mode 100644 index 00000000000..76d82d4184e --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_x.cl @@ -0,0 +1,23 @@ +/* + * Copyright 2019 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "kernel/kernel_compat_opencl.h" +#include "kernel/split/kernel_split_common.h" +#include "kernel/split/kernel_adaptive_filter_x.h" + +#define KERNEL_NAME adaptive_filter_x +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_y.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_y.cl new file mode 100644 index 00000000000..1e6d15ba0f2 --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_filter_y.cl @@ -0,0 +1,23 @@ +/* + * Copyright 2019 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "kernel/kernel_compat_opencl.h" +#include "kernel/split/kernel_split_common.h" +#include "kernel/split/kernel_adaptive_filter_y.h" + +#define KERNEL_NAME adaptive_filter_y +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_adaptive_stopping.cl b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_stopping.cl new file mode 100644 index 00000000000..51de0059667 --- /dev/null +++ b/intern/cycles/kernel/kernels/opencl/kernel_adaptive_stopping.cl @@ -0,0 +1,23 @@ +/* + * Copyright 2019 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "kernel/kernel_compat_opencl.h" +#include "kernel/split/kernel_split_common.h" +#include "kernel/split/kernel_adaptive_stopping.h" + +#define KERNEL_NAME adaptive_stopping +#include "kernel/kernels/opencl/kernel_split_function.h" +#undef KERNEL_NAME diff --git a/intern/cycles/kernel/kernels/opencl/kernel_split_bundle.cl b/intern/cycles/kernel/kernels/opencl/kernel_split_bundle.cl index 6041f13b52b..c3b7b09460a 100644 --- a/intern/cycles/kernel/kernels/opencl/kernel_split_bundle.cl +++ b/intern/cycles/kernel/kernels/opencl/kernel_split_bundle.cl @@ -28,3 +28,7 @@ #include "kernel/kernels/opencl/kernel_next_iteration_setup.cl" #include "kernel/kernels/opencl/kernel_indirect_subsurface.cl" #include "kernel/kernels/opencl/kernel_buffer_update.cl" +#include "kernel/kernels/opencl/kernel_adaptive_stopping.cl" +#include "kernel/kernels/opencl/kernel_adaptive_filter_x.cl" +#include "kernel/kernels/opencl/kernel_adaptive_filter_y.cl" +#include "kernel/kernels/opencl/kernel_adaptive_adjust_samples.cl" |