diff options
| author | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2017-09-27 02:03:50 +0300 |
|---|---|---|
| committer | Brecht Van Lommel <brechtvanlommel@gmail.com> | 2017-10-04 22:11:14 +0300 |
| commit | 12f453820514e9478afdda0acf4c4fb1eac11e1c (patch) | |
| tree | e8f9293c814457361febf1908e7131b0dc9ddfbd /intern/cycles/kernel/kernels | |
| parent | e3e16cecc4f080edbbd14e4bf1cfc580c5957d62 (diff) | |
Code refactor: use split variance calculation for mega kernels too.
There is no significant difference in denoised benchmark scenes and
denoising ctests, so might as well make it all consistent.
Diffstat (limited to 'intern/cycles/kernel/kernels')
| -rw-r--r-- | intern/cycles/kernel/kernels/cpu/filter_cpu.h | 6 | ||||
| -rw-r--r-- | intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h | 12 | ||||
| -rw-r--r-- | intern/cycles/kernel/kernels/cuda/filter.cu | 12 | ||||
| -rw-r--r-- | intern/cycles/kernel/kernels/opencl/filter.cl | 12 |
4 files changed, 14 insertions, 28 deletions
diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu.h b/intern/cycles/kernel/kernels/cpu/filter_cpu.h index 2ed713299fd..bf13ba62806 100644 --- a/intern/cycles/kernel/kernels/cpu/filter_cpu.h +++ b/intern/cycles/kernel/kernels/cpu/filter_cpu.h @@ -27,8 +27,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample, float *bufferV, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance); + int buffer_denoising_offset); void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, TilesInfo *tiles, @@ -40,8 +39,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, float *variance, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance); + int buffer_denoising_offset); void KERNEL_FUNCTION_FULL_NAME(filter_detect_outliers)(int x, int y, ccl_global float *image, diff --git a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h index 8dc1a8d583c..2fbb0ea2bdb 100644 --- a/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h +++ b/intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h @@ -45,8 +45,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample, float *bufferVariance, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_divide_shadow); @@ -60,8 +59,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_divide_shadow)(int sample, bufferVariance, load_int4(prefilter_rect), buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); #endif } @@ -74,8 +72,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, float *mean, float *variance, int* prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { #ifdef KERNEL_STUB STUB_ASSERT(KERNEL_ARCH, filter_get_feature); @@ -86,8 +83,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample, mean, variance, load_int4(prefilter_rect), buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); #endif } diff --git a/intern/cycles/kernel/kernels/cuda/filter.cu b/intern/cycles/kernel/kernels/cuda/filter.cu index 009c3fde9d5..c8172355a7f 100644 --- a/intern/cycles/kernel/kernels/cuda/filter.cu +++ b/intern/cycles/kernel/kernels/cuda/filter.cu @@ -37,8 +37,7 @@ kernel_cuda_filter_divide_shadow(int sample, float *bufferVariance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x; int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y; @@ -53,8 +52,7 @@ kernel_cuda_filter_divide_shadow(int sample, bufferVariance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } @@ -68,8 +66,7 @@ kernel_cuda_filter_get_feature(int sample, float *variance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - bool use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + blockDim.x*blockIdx.x + threadIdx.x; int y = prefilter_rect.y + blockDim.y*blockIdx.y + threadIdx.y; @@ -81,8 +78,7 @@ kernel_cuda_filter_get_feature(int sample, mean, variance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } diff --git a/intern/cycles/kernel/kernels/opencl/filter.cl b/intern/cycles/kernel/kernels/opencl/filter.cl index f015ac47d8a..7a7b596a350 100644 --- a/intern/cycles/kernel/kernels/opencl/filter.cl +++ b/intern/cycles/kernel/kernels/opencl/filter.cl @@ -31,8 +31,7 @@ __kernel void kernel_ocl_filter_divide_shadow(int sample, ccl_global float *bufferVariance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - char use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + get_global_id(0); int y = prefilter_rect.y + get_global_id(1); @@ -47,8 +46,7 @@ __kernel void kernel_ocl_filter_divide_shadow(int sample, bufferVariance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } @@ -60,8 +58,7 @@ __kernel void kernel_ocl_filter_get_feature(int sample, ccl_global float *variance, int4 prefilter_rect, int buffer_pass_stride, - int buffer_denoising_offset, - char use_split_variance) + int buffer_denoising_offset) { int x = prefilter_rect.x + get_global_id(0); int y = prefilter_rect.y + get_global_id(1); @@ -73,8 +70,7 @@ __kernel void kernel_ocl_filter_get_feature(int sample, mean, variance, prefilter_rect, buffer_pass_stride, - buffer_denoising_offset, - use_split_variance); + buffer_denoising_offset); } } |