Merge branch 'blender2.7'
authorBrecht Van Lommel <brechtvanlommel@gmail.com>
Wed, 6 Feb 2019 14:22:53 +0000 (15:22 +0100)
committerBrecht Van Lommel <brechtvanlommel@gmail.com>
Wed, 6 Feb 2019 14:22:53 +0000 (15:22 +0100)
31 files changed:
intern/cycles/blender/addon/engine.py
intern/cycles/blender/blender_session.cpp
intern/cycles/blender/blender_sync.cpp
intern/cycles/device/device_cpu.cpp
intern/cycles/device/device_cuda.cpp
intern/cycles/device/device_denoising.cpp
intern/cycles/device/device_denoising.h
intern/cycles/device/device_task.h
intern/cycles/device/opencl/opencl.h
intern/cycles/device/opencl/opencl_base.cpp
intern/cycles/kernel/filter/filter_defines.h
intern/cycles/kernel/filter/filter_features.h
intern/cycles/kernel/filter/filter_features_sse.h
intern/cycles/kernel/filter/filter_nlm_cpu.h
intern/cycles/kernel/filter/filter_nlm_gpu.h
intern/cycles/kernel/filter/filter_prefilter.h
intern/cycles/kernel/filter/filter_reconstruction.h
intern/cycles/kernel/filter/filter_transform.h
intern/cycles/kernel/filter/filter_transform_gpu.h
intern/cycles/kernel/filter/filter_transform_sse.h
intern/cycles/kernel/kernel_types.h
intern/cycles/kernel/kernels/cpu/filter_cpu.h
intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h
intern/cycles/kernel/kernels/cuda/filter.cu
intern/cycles/kernel/kernels/opencl/filter.cl
intern/cycles/render/buffers.cpp
intern/cycles/render/buffers.h
intern/cycles/render/film.cpp
intern/cycles/render/film.h
intern/cycles/render/session.cpp
intern/cycles/render/session.h

index 7829e09..b8bc74f 100644 (file)
@@ -270,14 +270,11 @@ def register_passes(engine, scene, srl):
         engine.register_pass(scene, srl, "Noisy Image", 4, "RGBA", 'COLOR')
         if crl.denoising_store_passes:
             engine.register_pass(scene, srl, "Denoising Normal",          3, "XYZ", 'VECTOR')
-            engine.register_pass(scene, srl, "Denoising Normal Variance", 3, "XYZ", 'VECTOR')
             engine.register_pass(scene, srl, "Denoising Albedo",          3, "RGB", 'COLOR')
-            engine.register_pass(scene, srl, "Denoising Albedo Variance", 3, "RGB", 'COLOR')
             engine.register_pass(scene, srl, "Denoising Depth",           1, "Z",   'VALUE')
-            engine.register_pass(scene, srl, "Denoising Depth Variance",  1, "Z",   'VALUE')
-            engine.register_pass(scene, srl, "Denoising Shadow A",        3, "XYV", 'VECTOR')
-            engine.register_pass(scene, srl, "Denoising Shadow B",        3, "XYV", 'VECTOR')
-            engine.register_pass(scene, srl, "Denoising Image Variance",  3, "RGB", 'COLOR')
+            engine.register_pass(scene, srl, "Denoising Shadowing",       1, "X",   'VALUE')
+            engine.register_pass(scene, srl, "Denoising Variance",        3, "RGB", 'COLOR')
+            engine.register_pass(scene, srl, "Denoising Intensity",       1, "X",   'VALUE')
             clean_options = ("denoising_diffuse_direct", "denoising_diffuse_indirect",
                              "denoising_glossy_direct", "denoising_glossy_indirect",
                              "denoising_transmission_direct", "denoising_transmission_indirect",
index 10170c9..6e6d98b 100644 (file)
@@ -423,15 +423,19 @@ void BlenderSession::render(BL::Depsgraph& b_depsgraph_)
        buffer_params.passes = passes;
 
        PointerRNA crl = RNA_pointer_get(&b_view_layer.ptr, "cycles");
-       bool use_denoising = get_boolean(crl, "use_denoising");
-       bool denoising_passes = use_denoising || get_boolean(crl, "denoising_store_passes");
+       bool full_denoising = get_boolean(crl, "use_denoising");
+       bool write_denoising_passes = get_boolean(crl, "denoising_store_passes");
 
-       session->tile_manager.schedule_denoising = use_denoising;
-       buffer_params.denoising_data_pass = denoising_passes;
+       bool run_denoising = full_denoising || write_denoising_passes;
+
+       session->tile_manager.schedule_denoising = run_denoising;
+       buffer_params.denoising_data_pass = run_denoising;
        buffer_params.denoising_clean_pass = (scene->film->denoising_flags & DENOISING_CLEAN_ALL_PASSES);
+       buffer_params.denoising_prefiltered_pass = write_denoising_passes;
 
-       session->params.use_denoising = use_denoising;
-       session->params.denoising_passes = denoising_passes;
+       session->params.run_denoising = run_denoising;
+       session->params.full_denoising = full_denoising;
+       session->params.write_denoising_passes = write_denoising_passes;
        session->params.denoising_radius = get_int(crl, "denoising_radius");
        session->params.denoising_strength = get_float(crl, "denoising_strength");
        session->params.denoising_feature_strength = get_float(crl, "denoising_feature_strength");
@@ -439,6 +443,7 @@ void BlenderSession::render(BL::Depsgraph& b_depsgraph_)
 
        scene->film->denoising_data_pass = buffer_params.denoising_data_pass;
        scene->film->denoising_clean_pass = buffer_params.denoising_clean_pass;
+       scene->film->denoising_prefiltered_pass = buffer_params.denoising_prefiltered_pass;
        session->params.denoising_radius = get_int(crl, "denoising_radius");
        session->params.denoising_strength = get_float(crl, "denoising_strength");
        session->params.denoising_feature_strength = get_float(crl, "denoising_feature_strength");
index e41a80a..072af28 100644 (file)
@@ -482,7 +482,7 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
 {
        string name = b_pass.name();
 
-       if(name == "Noisy Image") return DENOISING_PASS_COLOR;
+       if(name == "Noisy Image") return DENOISING_PASS_PREFILTERED_COLOR;
 
        if(name.substr(0, 10) != "Denoising ") {
                return -1;
@@ -490,15 +490,12 @@ int BlenderSync::get_denoising_pass(BL::RenderPass& b_pass)
        name = name.substr(10);
 
 #define MAP_PASS(passname, offset) if(name == passname) return offset;
-       MAP_PASS("Normal", DENOISING_PASS_NORMAL);
-       MAP_PASS("Normal Variance", DENOISING_PASS_NORMAL_VAR);
-       MAP_PASS("Albedo", DENOISING_PASS_ALBEDO);
-       MAP_PASS("Albedo Variance", DENOISING_PASS_ALBEDO_VAR);
-       MAP_PASS("Depth", DENOISING_PASS_DEPTH);
-       MAP_PASS("Depth Variance", DENOISING_PASS_DEPTH_VAR);
-       MAP_PASS("Shadow A", DENOISING_PASS_SHADOW_A);
-       MAP_PASS("Shadow B", DENOISING_PASS_SHADOW_B);
-       MAP_PASS("Image Variance", DENOISING_PASS_COLOR_VAR);
+       MAP_PASS("Normal", DENOISING_PASS_PREFILTERED_NORMAL);
+       MAP_PASS("Albedo", DENOISING_PASS_PREFILTERED_ALBEDO);
+       MAP_PASS("Depth", DENOISING_PASS_PREFILTERED_DEPTH);
+       MAP_PASS("Shadowing", DENOISING_PASS_PREFILTERED_SHADOWING);
+       MAP_PASS("Variance", DENOISING_PASS_PREFILTERED_VARIANCE);
+       MAP_PASS("Intensity", DENOISING_PASS_PREFILTERED_INTENSITY);
        MAP_PASS("Clean", DENOISING_PASS_CLEAN);
 #undef MAP_PASS
 
@@ -530,10 +527,11 @@ vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
        }
 
        PointerRNA crp = RNA_pointer_get(&b_view_layer.ptr, "cycles");
-       bool use_denoising = get_boolean(crp, "use_denoising");
-       bool store_denoising_passes = get_boolean(crp, "denoising_store_passes");
+       bool full_denoising = get_boolean(crp, "use_denoising");
+       bool write_denoising_passes = get_boolean(crp, "denoising_store_passes");
+
        scene->film->denoising_flags = 0;
-       if(use_denoising || store_denoising_passes) {
+       if(full_denoising || write_denoising_passes) {
 #define MAP_OPTION(name, flag) if(!get_boolean(crp, name)) scene->film->denoising_flags |= flag;
                MAP_OPTION("denoising_diffuse_direct",        DENOISING_CLEAN_DIFFUSE_DIR);
                MAP_OPTION("denoising_diffuse_indirect",      DENOISING_CLEAN_DIFFUSE_IND);
@@ -547,16 +545,13 @@ vector<Pass> BlenderSync::sync_render_passes(BL::RenderLayer& b_rlay,
                b_engine.add_pass("Noisy Image", 4, "RGBA", b_view_layer.name().c_str());
        }
 
-       if(store_denoising_passes) {
+       if(write_denoising_passes) {
                b_engine.add_pass("Denoising Normal",          3, "XYZ", b_view_layer.name().c_str());
-               b_engine.add_pass("Denoising Normal Variance", 3, "XYZ", b_view_layer.name().c_str());
                b_engine.add_pass("Denoising Albedo",          3, "RGB", b_view_layer.name().c_str());
-               b_engine.add_pass("Denoising Albedo Variance", 3, "RGB", b_view_layer.name().c_str());
                b_engine.add_pass("Denoising Depth",           1, "Z",   b_view_layer.name().c_str());
-               b_engine.add_pass("Denoising Depth Variance",  1, "Z",   b_view_layer.name().c_str());
-               b_engine.add_pass("Denoising Shadow A",        3, "XYV", b_view_layer.name().c_str());
-               b_engine.add_pass("Denoising Shadow B",        3, "XYV", b_view_layer.name().c_str());
-               b_engine.add_pass("Denoising Image Variance",  3, "RGB", b_view_layer.name().c_str());
+               b_engine.add_pass("Denoising Shadowing",       1, "X",   b_view_layer.name().c_str());
+               b_engine.add_pass("Denoising Variance",        3, "RGB", b_view_layer.name().c_str());
+               b_engine.add_pass("Denoising Intensity",       1, "X",   b_view_layer.name().c_str());
 
                if(scene->film->denoising_flags & DENOISING_CLEAN_ALL_PASSES) {
                        b_engine.add_pass("Denoising Clean",   3, "RGB", b_view_layer.name().c_str());
index a92c052..1f39a41 100644 (file)
@@ -180,20 +180,21 @@ public:
        KernelFunctions<void(*)(KernelGlobals *, uchar4 *, float *, float, int, int, int, int)> convert_to_byte_kernel;
        KernelFunctions<void(*)(KernelGlobals *, uint4 *, float4 *, int, int, int, int, int)>   shader_kernel;
 
-       KernelFunctions<void(*)(int, TileInfo*, int, int, float*, float*, float*, float*, float*, int*, int, int)> filter_divide_shadow_kernel;
-       KernelFunctions<void(*)(int, TileInfo*, int, int, int, int, float*, float*, int*, int, int)>               filter_get_feature_kernel;
+       KernelFunctions<void(*)(int, TileInfo*, int, int, float*, float*, float*, float*, float*, int*, int, int)>  filter_divide_shadow_kernel;
+       KernelFunctions<void(*)(int, TileInfo*, int, int, int, int, float*, float*, float, int*, int, int)>         filter_get_feature_kernel;
+       KernelFunctions<void(*)(int, int, int, int*, float*, float*, int, int*)>                                    filter_write_feature_kernel;
        KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)>                               filter_detect_outliers_kernel;
        KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int)>                               filter_combine_halves_kernel;
 
-       KernelFunctions<void(*)(int, int, float*, float*, float*, int*, int, int, float, float)>   filter_nlm_calc_difference_kernel;
-       KernelFunctions<void(*)(float*, float*, int*, int, int)>                                   filter_nlm_blur_kernel;
-       KernelFunctions<void(*)(float*, float*, int*, int, int)>                                   filter_nlm_calc_weight_kernel;
-       KernelFunctions<void(*)(int, int, float*, float*, float*, float*, float*, int*, int, int)> filter_nlm_update_output_kernel;
-       KernelFunctions<void(*)(float*, float*, int*, int)>                                        filter_nlm_normalize_kernel;
+       KernelFunctions<void(*)(int, int, float*, float*, float*, float*, int*, int, int, int, float, float)> filter_nlm_calc_difference_kernel;
+       KernelFunctions<void(*)(float*, float*, int*, int, int)>                                              filter_nlm_blur_kernel;
+       KernelFunctions<void(*)(float*, float*, int*, int, int)>                                              filter_nlm_calc_weight_kernel;
+       KernelFunctions<void(*)(int, int, float*, float*, float*, float*, float*, int*, int, int, int)>       filter_nlm_update_output_kernel;
+       KernelFunctions<void(*)(float*, float*, int*, int)>                                                   filter_nlm_normalize_kernel;
 
-       KernelFunctions<void(*)(float*, int, int, int, float*, int*, int*, int, int, float)>                         filter_construct_transform_kernel;
-       KernelFunctions<void(*)(int, int, float*, float*, float*, int*, float*, float3*, int*, int*, int, int, int)> filter_nlm_construct_gramian_kernel;
-       KernelFunctions<void(*)(int, int, int, float*, int*, float*, float3*, int*, int)>                            filter_finalize_kernel;
+       KernelFunctions<void(*)(float*, TileInfo*, int, int, int, float*, int*, int*, int, int, bool, int, float)>                   filter_construct_transform_kernel;
+       KernelFunctions<void(*)(int, int, int, float*, float*, float*, int*, float*, float3*, int*, int*, int, int, int, int, bool)> filter_nlm_construct_gramian_kernel;
+       KernelFunctions<void(*)(int, int, int, float*, int*, float*, float3*, int*, int)>                                            filter_finalize_kernel;
 
        KernelFunctions<void(*)(KernelGlobals *, ccl_constant KernelData*, ccl_global void*, int, ccl_global char*,
                               int, int, int, int, int, int, int, int, ccl_global int*, int,
@@ -218,6 +219,7 @@ public:
          REGISTER_KERNEL(shader),
          REGISTER_KERNEL(filter_divide_shadow),
          REGISTER_KERNEL(filter_get_feature),
+         REGISTER_KERNEL(filter_write_feature),
          REGISTER_KERNEL(filter_detect_outliers),
          REGISTER_KERNEL(filter_combine_halves),
          REGISTER_KERNEL(filter_nlm_calc_difference),
@@ -487,6 +489,8 @@ public:
 
                int w = align_up(rect.z-rect.x, 4);
                int h = rect.w-rect.y;
+               int stride = task->buffer.stride;
+               int channel_offset = task->nlm_state.is_color? task->buffer.pass_stride : 0;
 
                float *temporary_mem = (float*) task->buffer.temporary_mem.device_pointer;
                float *blurDifference = temporary_mem;
@@ -504,10 +508,11 @@ public:
                        filter_nlm_calc_difference_kernel()(dx, dy,
                                                            (float*) guide_ptr,
                                                            (float*) variance_ptr,
+                                                           NULL,
                                                            difference,
                                                            local_rect,
-                                                           w, 0,
-                                                           a, k_2);
+                                                           w, channel_offset,
+                                                           0, a, k_2);
 
                        filter_nlm_blur_kernel()       (difference, blurDifference, local_rect, w, f);
                        filter_nlm_calc_weight_kernel()(blurDifference, difference, local_rect, w, f);
@@ -520,7 +525,8 @@ public:
                                                          (float*) out_ptr,
                                                          weightAccum,
                                                          local_rect,
-                                                         w, f);
+                                                         channel_offset,
+                                                         stride, f);
                }
 
                int local_rect[4] = {0, 0, rect.z-rect.x, rect.w-rect.y};
@@ -536,6 +542,7 @@ public:
                for(int y = 0; y < task->filter_area.w; y++) {
                        for(int x = 0; x < task->filter_area.z; x++) {
                                filter_construct_transform_kernel()((float*) task->buffer.mem.device_pointer,
+                                                                   task->tile_info,
                                                                    x + task->filter_area.x,
                                                                    y + task->filter_area.y,
                                                                    y*task->filter_area.z + x,
@@ -543,6 +550,8 @@ public:
                                                                    (int*)   task->storage.rank.device_pointer,
                                                                    &task->rect.x,
                                                                    task->buffer.pass_stride,
+                                                                   task->buffer.frame_stride,
+                                                                   task->buffer.use_time,
                                                                    task->radius,
                                                                    task->pca_threshold);
                        }
@@ -550,21 +559,20 @@ public:
                return true;
        }
 
-       bool denoising_reconstruct(device_ptr color_ptr,
-                                  device_ptr color_variance_ptr,
-                                  device_ptr output_ptr,
-                                  DenoisingTask *task)
+       bool denoising_accumulate(device_ptr color_ptr,
+                                 device_ptr color_variance_ptr,
+                                 device_ptr scale_ptr,
+                                 int frame,
+                                 DenoisingTask *task)
        {
                ProfilingHelper profiling(task->profiler, PROFILING_DENOISING_RECONSTRUCT);
 
-               mem_zero(task->storage.XtWX);
-               mem_zero(task->storage.XtWY);
-
                float *temporary_mem = (float*) task->buffer.temporary_mem.device_pointer;
                float *difference     = temporary_mem;
                float *blurDifference = temporary_mem + task->buffer.pass_stride;
 
                int r = task->radius;
+               int frame_offset = frame * task->buffer.frame_stride;
                for(int i = 0; i < (2*r+1)*(2*r+1); i++) {
                        int dy = i / (2*r+1) - r;
                        int dx = i % (2*r+1) - r;
@@ -575,16 +583,19 @@ public:
                        filter_nlm_calc_difference_kernel()(dx, dy,
                                                            (float*) color_ptr,
                                                            (float*) color_variance_ptr,
+                                                           (float*) scale_ptr,
                                                            difference,
                                                            local_rect,
                                                            task->buffer.stride,
                                                            task->buffer.pass_stride,
+                                                           frame_offset,
                                                            1.0f,
                                                            task->nlm_k_2);
                        filter_nlm_blur_kernel()(difference, blurDifference, local_rect, task->buffer.stride, 4);
                        filter_nlm_calc_weight_kernel()(blurDifference, difference, local_rect, task->buffer.stride, 4);
                        filter_nlm_blur_kernel()(difference, blurDifference, local_rect, task->buffer.stride, 4);
                        filter_nlm_construct_gramian_kernel()(dx, dy,
+                                                             task->tile_info->frames[frame],
                                                              blurDifference,
                                                              (float*)  task->buffer.mem.device_pointer,
                                                              (float*)  task->storage.transform.device_pointer,
@@ -595,8 +606,17 @@ public:
                                                              &task->reconstruction_state.filter_window.x,
                                                              task->buffer.stride,
                                                              4,
-                                                             task->buffer.pass_stride);
+                                                             task->buffer.pass_stride,
+                                                             frame_offset,
+                                                             task->buffer.use_time);
                }
+
+               return true;
+       }
+
+       bool denoising_solve(device_ptr output_ptr,
+                            DenoisingTask *task)
+       {
                for(int y = 0; y < task->filter_area.w; y++) {
                        for(int x = 0; x < task->filter_area.z; x++) {
                                filter_finalize_kernel()(x,
@@ -661,6 +681,7 @@ public:
                                   int variance_offset,
                                   device_ptr mean_ptr,
                                   device_ptr variance_ptr,
+                                  float scale,
                                   DenoisingTask *task)
        {
                ProfilingHelper profiling(task->profiler, PROFILING_DENOISING_GET_FEATURE);
@@ -674,6 +695,7 @@ public:
                                                            x, y,
                                                            (float*) mean_ptr,
                                                            (float*) variance_ptr,
+                                                           scale,
                                                            &task->rect.x,
                                                            task->render_buffer.pass_stride,
                                                            task->render_buffer.offset);
@@ -682,6 +704,26 @@ public:
                return true;
        }
 
+       bool denoising_write_feature(int out_offset,
+                                    device_ptr from_ptr,
+                                    device_ptr buffer_ptr,
+                                    DenoisingTask *task)
+       {
+               for(int y = 0; y < task->filter_area.w; y++) {
+                       for(int x = 0; x < task->filter_area.z; x++) {
+                               filter_write_feature_kernel()(task->render_buffer.samples,
+                                                             x + task->filter_area.x,
+                                                             y + task->filter_area.y,
+                                                             &task->reconstruction_state.buffer_params.x,
+                                                             (float*) from_ptr,
+                                                             (float*) buffer_ptr,
+                                                             out_offset,
+                                                             &task->rect.x);
+                       }
+               }
+               return true;
+       }
+
        bool denoising_detect_outliers(device_ptr image_ptr,
                                       device_ptr variance_ptr,
                                       device_ptr depth_ptr,
@@ -754,11 +796,13 @@ public:
                tile.sample = tile.start_sample + tile.num_samples;
 
                denoising.functions.construct_transform = function_bind(&CPUDevice::denoising_construct_transform, this, &denoising);
-               denoising.functions.reconstruct = function_bind(&CPUDevice::denoising_reconstruct, this, _1, _2, _3, &denoising);
+               denoising.functions.accumulate = function_bind(&CPUDevice::denoising_accumulate, this, _1, _2, _3, _4, &denoising);
+               denoising.functions.solve = function_bind(&CPUDevice::denoising_solve, this, _1, &denoising);
                denoising.functions.divide_shadow = function_bind(&CPUDevice::denoising_divide_shadow, this, _1, _2, _3, _4, _5, &denoising);
                denoising.functions.non_local_means = function_bind(&CPUDevice::denoising_non_local_means, this, _1, _2, _3, _4, &denoising);
                denoising.functions.combine_halves = function_bind(&CPUDevice::denoising_combine_halves, this, _1, _2, _3, _4, _5, _6, &denoising);
-               denoising.functions.get_feature = function_bind(&CPUDevice::denoising_get_feature, this, _1, _2, _3, _4, &denoising);
+               denoising.functions.get_feature = function_bind(&CPUDevice::denoising_get_feature, this, _1, _2, _3, _4, _5, &denoising);
+               denoising.functions.write_feature = function_bind(&CPUDevice::denoising_write_feature, this, _1, _2, _3, &denoising);
                denoising.functions.detect_outliers = function_bind(&CPUDevice::denoising_detect_outliers, this, _1, _2, _3, _4, &denoising);
 
                denoising.filter_area = make_int4(tile.x, tile.y, tile.w, tile.h);
index 67f5793..ada538a 100644 (file)
@@ -1300,7 +1300,8 @@ public:
 
                int pass_stride = task->buffer.pass_stride;
                int num_shifts = (2*r+1)*(2*r+1);
-               int channel_offset = 0;
+               int channel_offset = task->nlm_state.is_color? task->buffer.pass_stride : 0;
+               int frame_offset = 0;
 
                if(have_error())
                        return false;
@@ -1308,6 +1309,7 @@ public:
                CUdeviceptr difference     = cuda_device_ptr(task->buffer.temporary_mem.device_pointer);
                CUdeviceptr blurDifference = difference + sizeof(float)*pass_stride*num_shifts;
                CUdeviceptr weightAccum = difference + 2*sizeof(float)*pass_stride*num_shifts;
+               CUdeviceptr scale_ptr = 0;
 
                cuda_assert(cuMemsetD8(weightAccum, 0, sizeof(float)*pass_stride));
                cuda_assert(cuMemsetD8(out_ptr, 0, sizeof(float)*pass_stride));
@@ -1326,10 +1328,10 @@ public:
 
                        CUDA_GET_BLOCKSIZE_1D(cuNLMCalcDifference, w*h, num_shifts);
 
-                       void *calc_difference_args[] = {&guide_ptr, &variance_ptr, &difference, &w, &h, &stride, &pass_stride, &r, &channel_offset, &a, &k_2};
+                       void *calc_difference_args[] = {&guide_ptr, &variance_ptr, &scale_ptr, &difference, &w, &h, &stride, &pass_stride, &r, &channel_offset, &frame_offset, &a, &k_2};
                        void *blur_args[]            = {&difference, &blurDifference, &w, &h, &stride, &pass_stride, &r, &f};
                        void *calc_weight_args[]     = {&blurDifference, &difference, &w, &h, &stride, &pass_stride, &r, &f};
-                       void *update_output_args[]   = {&blurDifference, &image_ptr, &out_ptr, &weightAccum, &w, &h, &stride, &pass_stride, &r, &f};
+                       void *update_output_args[]   = {&blurDifference, &image_ptr, &out_ptr, &weightAccum, &w, &h, &stride, &pass_stride, &channel_offset, &r, &f};
 
                        CUDA_LAUNCH_KERNEL_1D(cuNLMCalcDifference, calc_difference_args);
                        CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
@@ -1366,32 +1368,33 @@ public:
                                   task->storage.h);
 
                void *args[] = {&task->buffer.mem.device_pointer,
+                               &task->tile_info_mem.device_pointer,
                                &task->storage.transform.device_pointer,
                                &task->storage.rank.device_pointer,
                                &task->filter_area,
                                &task->rect,
                                &task->radius,
                                &task->pca_threshold,
-                               &task->buffer.pass_stride};
+                               &task->buffer.pass_stride,
+                               &task->buffer.frame_stride,
+                               &task->buffer.use_time};
                CUDA_LAUNCH_KERNEL(cuFilterConstructTransform, args);
                cuda_assert(cuCtxSynchronize());
 
                return !have_error();
        }
 
-       bool denoising_reconstruct(device_ptr color_ptr,
-                                  device_ptr color_variance_ptr,
-                                  device_ptr output_ptr,
-                                  DenoisingTask *task)
+       bool denoising_accumulate(device_ptr color_ptr,
+                                 device_ptr color_variance_ptr,
+                                 device_ptr scale_ptr,
+                                 int frame,
+                                 DenoisingTask *task)
        {
                if(have_error())
                        return false;
 
                CUDAContextScope scope(this);
 
-               mem_zero(task->storage.XtWX);
-               mem_zero(task->storage.XtWY);
-
                int r = task->radius;
                int f = 4;
                float a = 1.0f;
@@ -1400,6 +1403,8 @@ public:
                int w = task->reconstruction_state.source_w;
                int h = task->reconstruction_state.source_h;
                int stride = task->buffer.stride;
+               int frame_offset = frame * task->buffer.frame_stride;
+               int t = task->tile_info->frames[frame];
 
                int pass_stride = task->buffer.pass_stride;
                int num_shifts = (2*r+1)*(2*r+1);
@@ -1410,60 +1415,73 @@ public:
                CUdeviceptr difference     = cuda_device_ptr(task->buffer.temporary_mem.device_pointer);
                CUdeviceptr blurDifference = difference + sizeof(float)*pass_stride*num_shifts;
 
-               {
-                       CUfunction cuNLMCalcDifference, cuNLMBlur, cuNLMCalcWeight, cuNLMConstructGramian;
-                       cuda_assert(cuModuleGetFunction(&cuNLMCalcDifference,   cuFilterModule, "kernel_cuda_filter_nlm_calc_difference"));
-                       cuda_assert(cuModuleGetFunction(&cuNLMBlur,             cuFilterModule, "kernel_cuda_filter_nlm_blur"));
-                       cuda_assert(cuModuleGetFunction(&cuNLMCalcWeight,       cuFilterModule, "kernel_cuda_filter_nlm_calc_weight"));
-                       cuda_assert(cuModuleGetFunction(&cuNLMConstructGramian, cuFilterModule, "kernel_cuda_filter_nlm_construct_gramian"));
-
-                       cuda_assert(cuFuncSetCacheConfig(cuNLMCalcDifference,   CU_FUNC_CACHE_PREFER_L1));
-                       cuda_assert(cuFuncSetCacheConfig(cuNLMBlur,             CU_FUNC_CACHE_PREFER_L1));
-                       cuda_assert(cuFuncSetCacheConfig(cuNLMCalcWeight,       CU_FUNC_CACHE_PREFER_L1));
-                       cuda_assert(cuFuncSetCacheConfig(cuNLMConstructGramian, CU_FUNC_CACHE_PREFER_SHARED));
-
-                       CUDA_GET_BLOCKSIZE_1D(cuNLMCalcDifference,
-                                            task->reconstruction_state.source_w * task->reconstruction_state.source_h,
-                                            num_shifts);
-
-                       void *calc_difference_args[] = {&color_ptr, &color_variance_ptr, &difference, &w, &h, &stride, &pass_stride, &r, &pass_stride, &a, &k_2};
-                       void *blur_args[]            = {&difference, &blurDifference, &w, &h, &stride, &pass_stride, &r, &f};
-                       void *calc_weight_args[]     = {&blurDifference, &difference, &w, &h, &stride, &pass_stride, &r, &f};
-                       void *construct_gramian_args[] = {&blurDifference,
-                                                         &task->buffer.mem.device_pointer,
-                                                         &task->storage.transform.device_pointer,
-                                                         &task->storage.rank.device_pointer,
-                                                         &task->storage.XtWX.device_pointer,
-                                                         &task->storage.XtWY.device_pointer,
-                                                         &task->reconstruction_state.filter_window,
-                                                         &w, &h, &stride,
-                                                         &pass_stride, &r,
-                                                         &f};
-
-                       CUDA_LAUNCH_KERNEL_1D(cuNLMCalcDifference, calc_difference_args);
-                       CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
-                       CUDA_LAUNCH_KERNEL_1D(cuNLMCalcWeight, calc_weight_args);
-                       CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
-                       CUDA_LAUNCH_KERNEL_1D(cuNLMConstructGramian, construct_gramian_args);
-               }
+               CUfunction cuNLMCalcDifference, cuNLMBlur, cuNLMCalcWeight, cuNLMConstructGramian;
+               cuda_assert(cuModuleGetFunction(&cuNLMCalcDifference,   cuFilterModule, "kernel_cuda_filter_nlm_calc_difference"));
+               cuda_assert(cuModuleGetFunction(&cuNLMBlur,             cuFilterModule, "kernel_cuda_filter_nlm_blur"));
+               cuda_assert(cuModuleGetFunction(&cuNLMCalcWeight,       cuFilterModule, "kernel_cuda_filter_nlm_calc_weight"));
+               cuda_assert(cuModuleGetFunction(&cuNLMConstructGramian, cuFilterModule, "kernel_cuda_filter_nlm_construct_gramian"));
+
+               cuda_assert(cuFuncSetCacheConfig(cuNLMCalcDifference,   CU_FUNC_CACHE_PREFER_L1));
+               cuda_assert(cuFuncSetCacheConfig(cuNLMBlur,             CU_FUNC_CACHE_PREFER_L1));
+               cuda_assert(cuFuncSetCacheConfig(cuNLMCalcWeight,       CU_FUNC_CACHE_PREFER_L1));
+               cuda_assert(cuFuncSetCacheConfig(cuNLMConstructGramian, CU_FUNC_CACHE_PREFER_SHARED));
+
+               CUDA_GET_BLOCKSIZE_1D(cuNLMCalcDifference,
+                                    task->reconstruction_state.source_w * task->reconstruction_state.source_h,
+                                    num_shifts);
+
+               void *calc_difference_args[] = {&color_ptr,
+                                               &color_variance_ptr,
+                                               &scale_ptr,
+                                               &difference,
+                                               &w, &h,
+                                               &stride, &pass_stride,
+                                               &r, &pass_stride,
+                                               &frame_offset,
+                                               &a, &k_2};
+               void *blur_args[]            = {&difference, &blurDifference, &w, &h, &stride, &pass_stride, &r, &f};
+               void *calc_weight_args[]     = {&blurDifference, &difference, &w, &h, &stride, &pass_stride, &r, &f};
+               void *construct_gramian_args[] = {&t,
+                                                 &blurDifference,
+                                                 &task->buffer.mem.device_pointer,
+                                                 &task->storage.transform.device_pointer,
+                                                 &task->storage.rank.device_pointer,
+                                                 &task->storage.XtWX.device_pointer,
+                                                 &task->storage.XtWY.device_pointer,
+                                                 &task->reconstruction_state.filter_window,
+                                                 &w, &h, &stride,
+                                                 &pass_stride, &r,
+                                                 &f,
+                                                 &frame_offset,
+                                                 &task->buffer.use_time};
+
+               CUDA_LAUNCH_KERNEL_1D(cuNLMCalcDifference, calc_difference_args);
+               CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
+               CUDA_LAUNCH_KERNEL_1D(cuNLMCalcWeight, calc_weight_args);
+               CUDA_LAUNCH_KERNEL_1D(cuNLMBlur, blur_args);
+               CUDA_LAUNCH_KERNEL_1D(cuNLMConstructGramian, construct_gramian_args);
+               cuda_assert(cuCtxSynchronize());
 
-               {
-                       CUfunction cuFinalize;
-                       cuda_assert(cuModuleGetFunction(&cuFinalize, cuFilterModule, "kernel_cuda_filter_finalize"));
-                       cuda_assert(cuFuncSetCacheConfig(cuFinalize, CU_FUNC_CACHE_PREFER_L1));
-                       void *finalize_args[] = {&output_ptr,
-                                                        &task->storage.rank.device_pointer,
-                                                        &task->storage.XtWX.device_pointer,
-                                                        &task->storage.XtWY.device_pointer,
-                                                        &task->filter_area,
-                                                        &task->reconstruction_state.buffer_params.x,
-                                                        &task->render_buffer.samples};
-                       CUDA_GET_BLOCKSIZE(cuFinalize,
-                                          task->reconstruction_state.source_w,
-                                          task->reconstruction_state.source_h);
-                       CUDA_LAUNCH_KERNEL(cuFinalize, finalize_args);
-               }
+               return !have_error();
+       }
 
+       bool denoising_solve(device_ptr output_ptr,
+                            DenoisingTask *task)
+       {
+               CUfunction cuFinalize;
+               cuda_assert(cuModuleGetFunction(&cuFinalize, cuFilterModule, "kernel_cuda_filter_finalize"));
+               cuda_assert(cuFuncSetCacheConfig(cuFinalize, CU_FUNC_CACHE_PREFER_L1));
+               void *finalize_args[] = {&output_ptr,
+                                        &task->storage.rank.device_pointer,
+                                        &task->storage.XtWX.device_pointer,
+                                        &task->storage.XtWY.device_pointer,
+                                        &task->filter_area,
+                                        &task->reconstruction_state.buffer_params.x,
+                                        &task->render_buffer.samples};
+               CUDA_GET_BLOCKSIZE(cuFinalize,
+                                  task->reconstruction_state.source_w,
+                                  task->reconstruction_state.source_h);
+               CUDA_LAUNCH_KERNEL(cuFinalize, finalize_args);
                cuda_assert(cuCtxSynchronize());
 
                return !have_error();
@@ -1533,6 +1551,7 @@ public:
                                   int variance_offset,
                                   device_ptr mean_ptr,
                                   device_ptr variance_ptr,
+                                  float scale,
                                   DenoisingTask *task)
        {
                if(have_error())
@@ -1553,6 +1572,7 @@ public:
                                &variance_offset,
                                &mean_ptr,
                                &variance_ptr,
+                               &scale,
                                &task->rect,
                                &task->render_buffer.pass_stride,
                                &task->render_buffer.offset};
@@ -1562,6 +1582,36 @@ public:
                return !have_error();
        }
 
+       bool denoising_write_feature(int out_offset,
+                                    device_ptr from_ptr,
+                                    device_ptr buffer_ptr,
+                                    DenoisingTask *task)
+       {
+               if(have_error())
+                       return false;
+
+               CUDAContextScope scope(this);
+
+               CUfunction cuFilterWriteFeature;
+               cuda_assert(cuModuleGetFunction(&cuFilterWriteFeature, cuFilterModule, "kernel_cuda_filter_write_feature"));
+               cuda_assert(cuFuncSetCacheConfig(cuFilterWriteFeature, CU_FUNC_CACHE_PREFER_L1));
+               CUDA_GET_BLOCKSIZE(cuFilterWriteFeature,
+                                  task->filter_area.z,
+                                  task->filter_area.w);
+
+               void *args[] = {&task->render_buffer.samples,
+                               &task->reconstruction_state.buffer_params,
+                               &task->filter_area,
+                               &from_ptr,
+                               &buffer_ptr,
+                               &out_offset,
+                               &task->rect};
+               CUDA_LAUNCH_KERNEL(cuFilterWriteFeature, args);
+               cuda_assert(cuCtxSynchronize());
+
+               return !have_error();
+       }
+
        bool denoising_detect_outliers(device_ptr image_ptr,
                                       device_ptr variance_ptr,
                                       device_ptr depth_ptr,
@@ -1596,11 +1646,13 @@ public:
        void denoise(RenderTile &rtile, DenoisingTask& denoising)
        {
                denoising.functions.construct_transform = function_bind(&CUDADevice::denoising_construct_transform, this, &denoising);
-               denoising.functions.reconstruct = function_bind(&CUDADevice::denoising_reconstruct, this, _1, _2, _3, &denoising);
+               denoising.functions.accumulate = function_bind(&CUDADevice::denoising_accumulate, this, _1, _2, _3, _4, &denoising);
+               denoising.functions.solve = function_bind(&CUDADevice::denoising_solve, this, _1, &denoising);
                denoising.functions.divide_shadow = function_bind(&CUDADevice::denoising_divide_shadow, this, _1, _2, _3, _4, _5, &denoising);
                denoising.functions.non_local_means = function_bind(&CUDADevice::denoising_non_local_means, this, _1, _2, _3, _4, &denoising);
                denoising.functions.combine_halves = function_bind(&CUDADevice::denoising_combine_halves, this, _1, _2, _3, _4, _5, _6, &denoising);
-               denoising.functions.get_feature = function_bind(&CUDADevice::denoising_get_feature, this, _1, _2, _3, _4, &denoising);
+               denoising.functions.get_feature = function_bind(&CUDADevice::denoising_get_feature, this, _1, _2, _3, _4, _5, &denoising);
+               denoising.functions.write_feature = function_bind(&CUDADevice::denoising_write_feature, this, _1, _2, _3, &denoising);
                denoising.functions.detect_outliers = function_bind(&CUDADevice::denoising_detect_outliers, this, _1, _2, _3, _4, &denoising);
 
                denoising.filter_area = make_int4(rtile.x, rtile.y, rtile.w, rtile.h);
index 433cbd3..61e0ba4 100644 (file)
@@ -36,14 +36,28 @@ DenoisingTask::DenoisingTask(Device *device, const DeviceTask &task)
                pca_threshold = powf(10.0f, lerp(-5.0f, 3.0f, task.denoising_feature_strength));
        }
 
+       render_buffer.frame_stride = task.frame_stride;
        render_buffer.pass_stride = task.pass_stride;
        render_buffer.offset = task.pass_denoising_data;
 
-       target_buffer.pass_stride = task.pass_stride;
+       target_buffer.pass_stride = task.target_pass_stride;
        target_buffer.denoising_clean_offset = task.pass_denoising_clean;
+       target_buffer.offset = 0;
 
        functions.map_neighbor_tiles = function_bind(task.map_neighbor_tiles, _1, device);
        functions.unmap_neighbor_tiles = function_bind(task.unmap_neighbor_tiles, _1, device);
+
+       tile_info = (TileInfo*) tile_info_mem.alloc(sizeof(TileInfo)/sizeof(int));
+       tile_info->from_render = task.denoising_from_render? 1 : 0;
+
+       tile_info->frames[0] = 0;
+       tile_info->num_frames = min(task.denoising_frames.size() + 1, DENOISE_MAX_FRAMES);
+       for(int i = 1; i < tile_info->num_frames; i++) {
+               tile_info->frames[i] = task.denoising_frames[i-1];
+       }
+
+       write_passes = task.denoising_write_passes;
+       do_filter = task.denoising_do_filter;
 }
 
 DenoisingTask::~DenoisingTask()
@@ -59,8 +73,6 @@ DenoisingTask::~DenoisingTask()
 
 void DenoisingTask::set_render_buffer(RenderTile *rtiles)
 {
-       tile_info = (TileInfo*) tile_info_mem.alloc(sizeof(TileInfo)/sizeof(int));
-
        for(int i = 0; i < 9; i++) {
                tile_info->offsets[i] = rtiles[i].offset;
                tile_info->strides[i] = rtiles[i].stride;
@@ -79,6 +91,13 @@ void DenoisingTask::set_render_buffer(RenderTile *rtiles)
        target_buffer.stride = rtiles[9].stride;
        target_buffer.ptr    = rtiles[9].buffer;
 
+       if(write_passes && rtiles[9].buffers) {
+               target_buffer.denoising_output_offset = rtiles[9].buffers->params.get_denoising_prefiltered_offset();
+       }
+       else {
+               target_buffer.denoising_output_offset = 0;
+       }
+
        tile_info_mem.copy_to_device();
 }
 
@@ -89,15 +108,18 @@ void DenoisingTask::setup_denoising_buffer()
        rect = rect_expand(rect, radius);
        rect = rect_clip(rect, make_int4(tile_info->x[0], tile_info->y[0], tile_info->x[3], tile_info->y[3]));
 
-       buffer.passes = 14;
+       buffer.use_intensity = write_passes || (tile_info->num_frames > 1);
+       buffer.passes = buffer.use_intensity? 15 : 14;
        buffer.width = rect.z - rect.x;
        buffer.stride = align_up(buffer.width, 4);
        buffer.h = rect.w - rect.y;
        int alignment_floats = divide_up(device->mem_sub_ptr_alignment(), sizeof(float));
        buffer.pass_stride = align_up(buffer.stride * buffer.h, alignment_floats);
+       buffer.frame_stride = buffer.pass_stride * buffer.passes;
        /* Pad the total size by four floats since the SIMD kernels might go a bit over the end. */
-       int mem_size = align_up(buffer.pass_stride * buffer.passes + 4, alignment_floats);
+       int mem_size = align_up(tile_info->num_frames * buffer.frame_stride + 4, alignment_floats);
        buffer.mem.alloc_to_device(mem_size, false);
+       buffer.use_time = (tile_info->num_frames > 1);
 
        /* CPUs process shifts sequentially while GPUs process them in parallel. */
        int num_layers;
@@ -129,14 +151,14 @@ void DenoisingTask::prefilter_shadowing()
        functions.divide_shadow(*unfiltered_a, *unfiltered_b, *sample_var, *sample_var_var, *buffer_var);
 
        /* Smooth the (generally pretty noisy) buffer variance using the spatial information from the sample variance. */
-       nlm_state.set_parameters(6, 3, 4.0f, 1.0f);
+       nlm_state.set_parameters(6, 3, 4.0f, 1.0f, false);
        functions.non_local_means(*buffer_var, *sample_var, *sample_var_var, *filtered_var);
 
        /* Reuse memory, the previous data isn't needed anymore. */
        device_ptr filtered_a = *buffer_var,
                   filtered_b = *sample_var;
        /* Use the smoothed variance to filter the two shadow half images using each other for weight calculation. */
-       nlm_state.set_parameters(5, 3, 1.0f, 0.25f);
+       nlm_state.set_parameters(5, 3, 1.0f, 0.25f, false);
        functions.non_local_means(*unfiltered_a, *unfiltered_b, *filtered_var, filtered_a);
        functions.non_local_means(*unfiltered_b, *unfiltered_a, *filtered_var, filtered_b);
 
@@ -147,7 +169,7 @@ void DenoisingTask::prefilter_shadowing()
        device_ptr final_a = *unfiltered_a,
                   final_b = *unfiltered_b;
        /* Use the residual variance for a second filter pass. */
-       nlm_state.set_parameters(4, 2, 1.0f, 0.5f);
+       nlm_state.set_parameters(4, 2, 1.0f, 0.5f, false);
        functions.non_local_means(filtered_a, filtered_b, residual_var, final_a);
        functions.non_local_means(filtered_b, filtered_a, residual_var, final_b);
 
@@ -167,9 +189,9 @@ void DenoisingTask::prefilter_features()
        for(int pass = 0; pass < 7; pass++) {
                device_sub_ptr feature_pass(buffer.mem, pass_to[pass]*buffer.pass_stride, buffer.pass_stride);
                /* Get the unfiltered pass and its variance from the RenderBuffers. */
-               functions.get_feature(mean_from[pass], variance_from[pass], *unfiltered, *variance);
+               functions.get_feature(mean_from[pass], variance_from[pass], *unfiltered, *variance, 1.0f / render_buffer.samples);
                /* Smooth the pass and store the result in the denoising buffers. */
-               nlm_state.set_parameters(2, 2, 1.0f, 0.25f);
+               nlm_state.set_parameters(2, 2, 1.0f, 0.25f, false);
                functions.non_local_means(*unfiltered, *unfiltered, *variance, *feature_pass);
        }
 }
@@ -188,13 +210,52 @@ void DenoisingTask::prefilter_color()
        for(int pass = 0; pass < num_color_passes; pass++) {
                device_sub_ptr color_pass(temporary_color, pass*buffer.pass_stride, buffer.pass_stride);
                device_sub_ptr color_var_pass(buffer.mem, variance_to[pass]*buffer.pass_stride, buffer.pass_stride);
-               functions.get_feature(mean_from[pass], variance_from[pass], *color_pass, *color_var_pass);
+               functions.get_feature(mean_from[pass], variance_from[pass], *color_pass, *color_var_pass, 1.0f / render_buffer.samples);
        }
 
        device_sub_ptr depth_pass    (buffer.mem,                                 0,   buffer.pass_stride);
        device_sub_ptr color_var_pass(buffer.mem, variance_to[0]*buffer.pass_stride, 3*buffer.pass_stride);
        device_sub_ptr output_pass   (buffer.mem,     mean_to[0]*buffer.pass_stride, 3*buffer.pass_stride);
        functions.detect_outliers(temporary_color.device_pointer, *color_var_pass, *depth_pass, *output_pass);
+
+       if(buffer.use_intensity) {
+               device_sub_ptr intensity_pass(buffer.mem, 14*buffer.pass_stride, buffer.pass_stride);
+               nlm_state.set_parameters(radius, 4, 2.0f, nlm_k_2*4.0f, true);
+               functions.non_local_means(*output_pass, *output_pass, *color_var_pass, *intensity_pass);
+       }
+}
+
+void DenoisingTask::load_buffer()
+{
+       device_ptr null_ptr = (device_ptr) 0;
+
+       int original_offset = render_buffer.offset;
+
+       int num_passes = buffer.use_intensity? 15 : 14;
+       for(int i = 0; i < tile_info->num_frames; i++) {
+               for(int pass = 0; pass < num_passes; pass++) {
+                       device_sub_ptr to_pass(buffer.mem, i*buffer.frame_stride + pass*buffer.pass_stride, buffer.pass_stride);
+                       bool is_variance = (pass >= 11) && (pass <= 13);
+                       functions.get_feature(pass, -1, *to_pass, null_ptr, is_variance? (1.0f / render_buffer.samples) : 1.0f);
+               }
+               render_buffer.offset += render_buffer.frame_stride;
+       }
+
+       render_buffer.offset = original_offset;
+}
+
+void DenoisingTask::write_buffer()
+{
+       reconstruction_state.buffer_params = make_int4(target_buffer.offset,
+                                                      target_buffer.stride,
+                                                      target_buffer.pass_stride,
+                                                      target_buffer.denoising_clean_offset);
+       int num_passes = buffer.use_intensity? 15 : 14;
+       for(int pass = 0; pass < num_passes; pass++) {
+               device_sub_ptr from_pass(buffer.mem, pass*buffer.pass_stride, buffer.pass_stride);
+               int out_offset = pass + target_buffer.denoising_output_offset;
+               functions.write_feature(out_offset, *from_pass, target_buffer.ptr);
+       }
 }
 
 void DenoisingTask::construct_transform()
@@ -212,6 +273,8 @@ void DenoisingTask::reconstruct()
 {
        storage.XtWX.alloc_to_device(storage.w*storage.h*XTWX_SIZE, false);
        storage.XtWY.alloc_to_device(storage.w*storage.h*XTWY_SIZE, false);
+       storage.XtWX.zero_to_device();
+       storage.XtWY.zero_to_device();
 
        reconstruction_state.filter_window = rect_from_shape(filter_area.x-rect.x, filter_area.y-rect.y, storage.w, storage.h);
        int tile_coordinate_offset = filter_area.y*target_buffer.stride + filter_area.x;
@@ -224,7 +287,18 @@ void DenoisingTask::reconstruct()
 
        device_sub_ptr color_ptr    (buffer.mem,  8*buffer.pass_stride, 3*buffer.pass_stride);
        device_sub_ptr color_var_ptr(buffer.mem, 11*buffer.pass_stride, 3*buffer.pass_stride);
-       functions.reconstruct(*color_ptr, *color_var_ptr, target_buffer.ptr);
+       for(int f = 0; f < tile_info->num_frames; f++) {
+               device_ptr scale_ptr = 0;
+               device_sub_ptr *scale_sub_ptr = NULL;
+               if(tile_info->frames[f] != 0 && (tile_info->num_frames > 1)) {
+                       scale_sub_ptr = new device_sub_ptr(buffer.mem, 14*buffer.pass_stride, buffer.pass_stride);
+                       scale_ptr = **scale_sub_ptr;
+               }
+
+               functions.accumulate(*color_ptr, *color_var_ptr, scale_ptr, f);
+               delete scale_sub_ptr;
+       }
+       functions.solve(target_buffer.ptr);
 }
 
 void DenoisingTask::run_denoising(RenderTile *tile)
@@ -236,12 +310,23 @@ void DenoisingTask::run_denoising(RenderTile *tile)
 
        setup_denoising_buffer();
 
-       prefilter_shadowing();
-       prefilter_features();
-       prefilter_color();
+       if(tile_info->from_render) {
+               prefilter_shadowing();
+               prefilter_features();
+               prefilter_color();
+       }
+       else {
+               load_buffer();
+       }
+
+       if(do_filter) {
+               construct_transform();
+               reconstruct();
+       }
 
-       construct_transform();
-       reconstruct();
+       if(write_passes) {
+               write_buffer();
+       }
 
        functions.unmap_neighbor_tiles(rtiles);
 }
index beae60c..5869aa0 100644 (file)
@@ -38,6 +38,7 @@ public:
        struct RenderBuffers {
                int offset;
                int pass_stride;
+               int frame_stride;
                int samples;
        } render_buffer;
 
@@ -47,6 +48,7 @@ public:
                int stride;
                int pass_stride;
                int denoising_clean_offset;
+               int denoising_output_offset;
                device_ptr ptr;
        } target_buffer;
 
@@ -58,6 +60,9 @@ public:
        int4 rect;
        int4 filter_area;
 
+       bool write_passes;
+       bool do_filter;
+
        struct DeviceFunctions {
                function<bool(device_ptr image_ptr,    /* Contains the values that are smoothed. */
                              device_ptr guide_ptr,    /* Contains the values that are used to calculate weights. */
@@ -66,8 +71,10 @@ public:
                              )> non_local_means;
                function<bool(device_ptr color_ptr,
                              device_ptr color_variance_ptr,
-                             device_ptr output_ptr
-                             )> reconstruct;
+                             device_ptr scale_ptr,
+                             int frame
+                             )> accumulate;
+               function<bool(device_ptr output_ptr)> solve;
                function<bool()> construct_transform;
 
                function<bool(device_ptr a_ptr,
@@ -86,13 +93,18 @@ public:
                function<bool(int mean_offset,
                              int variance_offset,
                              device_ptr mean_ptr,
-                             device_ptr variance_ptr
+                             device_ptr variance_ptr,
+                             float scale
                              )> get_feature;
                function<bool(device_ptr image_ptr,
                              device_ptr variance_ptr,
                              device_ptr depth_ptr,
                              device_ptr output_ptr
                              )> detect_outliers;
+               function<bool(int out_offset,
+                             device_ptr frop_ptr,
+                             device_ptr buffer_ptr
+                             )> write_feature;
                function<void(RenderTile *rtiles)> map_neighbor_tiles;
                function<void(RenderTile *rtiles)> unmap_neighbor_tiles;
        } functions;
@@ -114,8 +126,9 @@ public:
                int f;      /* Patch size of the filter. */
                float a;    /* Variance compensation factor in the MSE estimation. */
                float k_2;  /* Squared value of the k parameter of the filter. */
+               bool is_color;
 
-               void set_parameters(int r_, int f_, float a_, float k_2_) { r = r_; f = f_; a = a_, k_2 = k_2_; }
+               void set_parameters(int r_, int f_, float a_, float k_2_, bool is_color_) { r = r_; f = f_; a = a_, k_2 = k_2_; is_color = is_color_; }
        } nlm_state;
 
        struct Storage {
@@ -145,8 +158,11 @@ public:
                int stride;
                int h;
                int width;
+               int frame_stride;
                device_only_memory<float> mem;
                device_only_memory<float> temporary_mem;
+               bool use_time;
+               bool use_intensity;
 
                bool gpu_temporary_mem;
 
@@ -166,6 +182,9 @@ protected:
        void prefilter_color();
        void construct_transform();
        void reconstruct();
+
+       void load_buffer();
+       void write_buffer();
 };
 
 CCL_NAMESPACE_END
index 8610143..2871bc5 100644 (file)
@@ -72,7 +72,15 @@ public:
        float denoising_strength;
        float denoising_feature_strength;
        bool denoising_relative_pca;
+       bool denoising_from_render;
+       vector<int> denoising_frames;
+
+       bool denoising_do_filter;
+       bool denoising_write_passes;
+
        int pass_stride;
+       int frame_stride;
+       int target_pass_stride;
        int pass_denoising_data;
        int pass_denoising_clean;
 
index ea7ed4f..9b76316 100644 (file)
@@ -419,10 +419,13 @@ protected:
                                       device_ptr out_ptr,
                                       DenoisingTask *task);
        bool denoising_construct_transform(DenoisingTask *task);
-       bool denoising_reconstruct(device_ptr color_ptr,
-                                  device_ptr color_variance_ptr,
-                                  device_ptr output_ptr,
-                                  DenoisingTask *task);
+       bool denoising_accumulate(device_ptr color_ptr,
+                                 device_ptr color_variance_ptr,
+                                 device_ptr scale_ptr,
+                                 int frame,
+                                 DenoisingTask *task);
+       bool denoising_solve(device_ptr output_ptr,
+                            DenoisingTask *task);
        bool denoising_combine_halves(device_ptr a_ptr,
                                      device_ptr b_ptr,
                                      device_ptr mean_ptr,
@@ -439,7 +442,12 @@ protected:
                                   int variance_offset,
                                   device_ptr mean_ptr,
                                   device_ptr variance_ptr,
+                                  float scale,
                                   DenoisingTask *task);
+       bool denoising_write_feature(int to_offset,
+                                    device_ptr from_ptr,
+                                    device_ptr buffer_ptr,
+                                    DenoisingTask *task);
        bool denoising_detect_outliers(device_ptr image_ptr,
                                       device_ptr variance_ptr,
                                       device_ptr depth_ptr,
index d4d7c0f..4417065 100644 (file)
@@ -748,6 +748,7 @@ bool OpenCLDeviceBase::denoising_non_local_means(device_ptr image_ptr,
 
        int pass_stride = task->buffer.pass_stride;
        int num_shifts = (2*r+1)*(2*r+1);
+       int channel_offset = task->nlm_state.is_color? task->buffer.pass_stride : 0;
 
        device_sub_ptr difference(task->buffer.temporary_mem, 0, pass_stride*num_shifts);
        device_sub_ptr blurDifference(task->buffer.temporary_mem, pass_stride*num_shifts, pass_stride*num_shifts);
@@ -760,6 +761,7 @@ bool OpenCLDeviceBase::denoising_non_local_means(device_ptr image_ptr,
        cl_mem guide_mem = CL_MEM_PTR(guide_ptr);
        cl_mem variance_mem = CL_MEM_PTR(variance_ptr);
        cl_mem out_mem = CL_MEM_PTR(out_ptr);
+       cl_mem scale_mem = NULL;
 
        mem_zero_kernel(*weightAccum, sizeof(float)*pass_stride);
        mem_zero_kernel(out_ptr, sizeof(float)*pass_stride);
@@ -773,10 +775,12 @@ bool OpenCLDeviceBase::denoising_non_local_means(device_ptr image_ptr,
        kernel_set_args(ckNLMCalcDifference, 0,
                        guide_mem,
                        variance_mem,
+                       scale_mem,
                        difference_mem,
                        w, h, stride,
                        pass_stride,
-                       r, 0, a, k_2);
+                       r, channel_offset,
+                       0, a, k_2);
        kernel_set_args(ckNLMBlur, 0,
                        difference_mem,
                        blurDifference_mem,
@@ -796,6 +800,7 @@ bool OpenCLDeviceBase::denoising_non_local_means(device_ptr image_ptr,
                        weightAccum_mem,
                        w, h, stride,
                        pass_stride,
+                       channel_offset,
                        r, f);
 
        enqueue_kernel(ckNLMCalcDifference, w*h, num_shifts, true);
@@ -816,16 +821,31 @@ bool OpenCLDeviceBase::denoising_construct_transform(DenoisingTask *task)
        cl_mem buffer_mem = CL_MEM_PTR(task->buffer.mem.device_pointer);
        cl_mem transform_mem = CL_MEM_PTR(task->storage.transform.device_pointer);
        cl_mem rank_mem = CL_MEM_PTR(task->storage.rank.device_pointer);
+       cl_mem tile_info_mem = CL_MEM_PTR(task->tile_info_mem.device_pointer);
+
+       char use_time = task->buffer.use_time? 1 : 0;
 
        cl_kernel ckFilterConstructTransform = denoising_program(ustring("filter_construct_transform"));
 
-       kernel_set_args(ckFilterConstructTransform, 0,
-                       buffer_mem,
+       int arg_ofs = kernel_set_args(ckFilterConstructTransform, 0,
+                                     buffer_mem,
+                                     tile_info_mem);
+       cl_mem buffers[9];
+       for(int i = 0; i < 9; i++) {
+               buffers[i] = CL_MEM_PTR(task->tile_info->buffers[i]);
+               arg_ofs += kernel_set_args(ckFilterConstructTransform,
+                                          arg_ofs,
+                                          buffers[i]);
+       }
+       kernel_set_args(ckFilterConstructTransform,
+                       arg_ofs,
                        transform_mem,
                        rank_mem,
                        task->filter_area,
                        task->rect,
                        task->buffer.pass_stride,
+                       task->buffer.frame_stride,
+                       use_time,
                        task->radius,
                        task->pca_threshold);
 
@@ -837,17 +857,15 @@ bool OpenCLDeviceBase::denoising_construct_transform(DenoisingTask *task)
        return true;
 }
 
-bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
-                                             device_ptr color_variance_ptr,
-                                             device_ptr output_ptr,
-                                             DenoisingTask *task)
+bool OpenCLDeviceBase::denoising_accumulate(device_ptr color_ptr,
+                                            device_ptr color_variance_ptr,
+                                            device_ptr scale_ptr,
+                                            int frame,
+                                            DenoisingTask *task)
 {
-       mem_zero(task->storage.XtWX);
-       mem_zero(task->storage.XtWY);
-
        cl_mem color_mem = CL_MEM_PTR(color_ptr);
        cl_mem color_variance_mem = CL_MEM_PTR(color_variance_ptr);
-       cl_mem output_mem = CL_MEM_PTR(output_ptr);
+       cl_mem scale_mem = CL_MEM_PTR(scale_ptr);
 
        cl_mem buffer_mem = CL_MEM_PTR(task->buffer.mem.device_pointer);
        cl_mem transform_mem = CL_MEM_PTR(task->storage.transform.device_pointer);
@@ -859,11 +877,13 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
        cl_kernel ckNLMBlur             = denoising_program(ustring("filter_nlm_blur"));
        cl_kernel ckNLMCalcWeight       = denoising_program(ustring("filter_nlm_calc_weight"));
        cl_kernel ckNLMConstructGramian = denoising_program(ustring("filter_nlm_construct_gramian"));
-       cl_kernel ckFinalize            = denoising_program(ustring("filter_finalize"));
 
        int w = task->reconstruction_state.source_w;
        int h = task->reconstruction_state.source_h;
        int stride = task->buffer.stride;
+       int frame_offset = frame * task->buffer.frame_stride;
+       int t = task->tile_info->frames[frame];
+       char use_time = task->buffer.use_time? 1 : 0;
 
        int r = task->radius;
        int pass_stride = task->buffer.pass_stride;
@@ -877,11 +897,13 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
        kernel_set_args(ckNLMCalcDifference, 0,
                        color_mem,
                        color_variance_mem,
+                       scale_mem,
                        difference_mem,
                        w, h, stride,
                        pass_stride,
                        r,
                        pass_stride,
+                       frame_offset,
                        1.0f, task->nlm_k_2);
        kernel_set_args(ckNLMBlur, 0,
                        difference_mem,
@@ -896,6 +918,7 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
                        pass_stride,
                        r, 4);
        kernel_set_args(ckNLMConstructGramian, 0,
+                       t,
                        blurDifference_mem,
                        buffer_mem,
                        transform_mem,
@@ -905,7 +928,9 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
                        task->reconstruction_state.filter_window,
                        w, h, stride,
                        pass_stride,
-                       r, 4);
+                       r, 4,
+                       frame_offset,
+                       use_time);
 
        enqueue_kernel(ckNLMCalcDifference,   w*h, num_shifts, true);
        enqueue_kernel(ckNLMBlur,             w*h, num_shifts, true);
@@ -913,6 +938,22 @@ bool OpenCLDeviceBase::denoising_reconstruct(device_ptr color_ptr,
        enqueue_kernel(ckNLMBlur,             w*h, num_shifts, true);
        enqueue_kernel(ckNLMConstructGramian, w*h, num_shifts, true, 256);
 
+       return true;
+}
+
+bool OpenCLDeviceBase::denoising_solve(device_ptr output_ptr,
+                                       DenoisingTask *task)
+{
+       cl_kernel ckFinalize = denoising_program(ustring("filter_finalize"));
+
+       cl_mem output_mem = CL_MEM_PTR(output_ptr);
+       cl_mem rank_mem   = CL_MEM_PTR(task->storage.rank.device_pointer);
+       cl_mem XtWX_mem   = CL_MEM_PTR(task->storage.XtWX.device_pointer);
+       cl_mem XtWY_mem   = CL_MEM_PTR(task->storage.XtWY.device_pointer);
+
+       int w = task->reconstruction_state.source_w;
+       int h = task->reconstruction_state.source_h;
+
        kernel_set_args(ckFinalize, 0,
                        output_mem,
                        rank_mem,
@@ -1000,6 +1041,7 @@ bool OpenCLDeviceBase::denoising_get_feature(int mean_offset,
                                              int variance_offset,
                                              device_ptr mean_ptr,
                                              device_ptr variance_ptr,
+                                             float scale,
                                              DenoisingTask *task)
 {
        cl_mem mean_mem = CL_MEM_PTR(mean_ptr);
@@ -1023,6 +1065,7 @@ bool OpenCLDeviceBase::denoising_get_feature(int mean_offset,
                        variance_offset,
                        mean_mem,
                        variance_mem,
+                       scale,
                        task->rect,
                        task->render_buffer.pass_stride,
                        task->render_buffer.offset);
@@ -1033,6 +1076,31 @@ bool OpenCLDeviceBase::denoising_get_feature(int mean_offset,
        return true;
 }
 
+bool OpenCLDeviceBase::denoising_write_feature(int out_offset,
+                                               device_ptr from_ptr,
+                                               device_ptr buffer_ptr,
+                                               DenoisingTask *task)
+{
+       cl_mem from_mem = CL_MEM_PTR(from_ptr);
+       cl_mem buffer_mem = CL_MEM_PTR(buffer_ptr);
+
+       cl_kernel ckFilterWriteFeature = denoising_program(ustring("filter_write_feature"));
+
+       kernel_set_args(ckFilterWriteFeature, 0,
+                       task->render_buffer.samples,
+                       task->reconstruction_state.buffer_params,
+                       task->filter_area,
+                       from_mem,
+                       buffer_mem,
+                       out_offset,
+                       task->rect);
+       enqueue_kernel(ckFilterWriteFeature,
+                      task->filter_area.z,
+                      task->filter_area.w);
+
+       return true;
+}
+
 bool OpenCLDeviceBase::denoising_detect_outliers(device_ptr image_ptr,
                                                  device_ptr variance_ptr,
                                                  device_ptr depth_ptr,
@@ -1063,11 +1131,13 @@ bool OpenCLDeviceBase::denoising_detect_outliers(device_ptr image_ptr,
 void OpenCLDeviceBase::denoise(RenderTile &rtile, DenoisingTask& denoising)
 {
        denoising.functions.construct_transform = function_bind(&OpenCLDeviceBase::denoising_construct_transform, this, &denoising);
-       denoising.functions.reconstruct = function_bind(&OpenCLDeviceBase::denoising_reconstruct, this, _1, _2, _3, &denoising);
+       denoising.functions.accumulate = function_bind(&OpenCLDeviceBase::denoising_accumulate, this, _1, _2, _3, _4, &denoising);
+       denoising.functions.solve = function_bind(&OpenCLDeviceBase::denoising_solve, this, _1, &denoising);
        denoising.functions.divide_shadow = function_bind(&OpenCLDeviceBase::denoising_divide_shadow, this, _1, _2, _3, _4, _5, &denoising);
        denoising.functions.non_local_means = function_bind(&OpenCLDeviceBase::denoising_non_local_means, this, _1, _2, _3, _4, &denoising);
        denoising.functions.combine_halves = function_bind(&OpenCLDeviceBase::denoising_combine_halves, this, _1, _2, _3, _4, _5, _6, &denoising);
-       denoising.functions.get_feature = function_bind(&OpenCLDeviceBase::denoising_get_feature, this, _1, _2, _3, _4, &denoising);
+       denoising.functions.get_feature = function_bind(&OpenCLDeviceBase::denoising_get_feature, this, _1, _2, _3, _4, _5, &denoising);
+       denoising.functions.write_feature = function_bind(&OpenCLDeviceBase::denoising_write_feature, this, _1, _2, _3, &denoising);
        denoising.functions.detect_outliers = function_bind(&OpenCLDeviceBase::denoising_detect_outliers, this, _1, _2, _3, _4, &denoising);
 
        denoising.filter_area = make_int4(rtile.x, rtile.y, rtile.w, rtile.h);
index 67f4e62..cb04aac 100644 (file)
 #ifndef __FILTER_DEFINES_H__
 #define __FILTER_DEFINES_H__
 
-#define DENOISE_FEATURES 10
+#define DENOISE_FEATURES 11
 #define TRANSFORM_SIZE (DENOISE_FEATURES*DENOISE_FEATURES)
 #define XTWX_SIZE      (((DENOISE_FEATURES+1)*(DENOISE_FEATURES+2))/2)
 #define XTWY_SIZE      (DENOISE_FEATURES+1)
 
+#define DENOISE_MAX_FRAMES 16
+
 typedef struct TileInfo {
        int offsets[9];
        int strides[9];
        int x[4];
        int y[4];
+       int from_render;
+       int frames[DENOISE_MAX_FRAMES];
+       int num_frames;
        /* TODO(lukas): CUDA doesn't have uint64_t... */
 #ifdef __KERNEL_OPENCL__
        ccl_global float *buffers[9];
index 6226ed2..e1ea648 100644 (file)
 
 #define ccl_get_feature(buffer, pass) (buffer)[(pass)*pass_stride]
 
-/* Loop over the pixels in the range [low.x, high.x) x [low.y, high.y).
- * pixel_buffer always points to the current pixel in the first pass. */
-#define FOR_PIXEL_WINDOW     pixel_buffer = buffer + (low.y - rect.y)*buffer_w + (low.x - rect.x); \
-                             for(pixel.y = low.y; pixel.y < high.y; pixel.y++) { \
-                                 for(pixel.x = low.x; pixel.x < high.x; pixel.x++, pixel_buffer++) {
+/* Loop over the pixels in the range [low.x, high.x) x [low.y, high.y).+ * pixel_buffer always points to the current pixel in the first pass.
+ * Repeat the loop for every secondary frame if there are any. */
+#define FOR_PIXEL_WINDOW     for(int frame = 0; frame < tile_info->num_frames; frame++) { \
+                                 pixel.z = tile_info->frames[frame]; \
+                                 pixel_buffer = buffer + (low.y - rect.y)*buffer_w + (low.x - rect.x) + frame*frame_stride; \
+                                 for(pixel.y = low.y; pixel.y < high.y; pixel.y++) { \
+                                     for(pixel.x = low.x; pixel.x < high.x; pixel.x++, pixel_buffer++) {
 
-#define END_FOR_PIXEL_WINDOW     } \
-                                 pixel_buffer += buffer_w - (high.x - low.x); \
+#define END_FOR_PIXEL_WINDOW         } \
+                                     pixel_buffer += buffer_w - (high.x - low.x); \
+                                 } \
                              }
 
-ccl_device_inline void filter_get_features(int2 pixel,
+ccl_device_inline void filter_get_features(int3 pixel,
                                            const ccl_global float *ccl_restrict buffer,
                                            float *features,
+                                           bool use_time,
                                            const float *ccl_restrict mean,
                                            int pass_stride)
 {
@@ -44,15 +48,20 @@ ccl_device_inline void filter_get_features(int2 pixel,
        features[7] = ccl_get_feature(buffer, 5);
        features[8] = ccl_get_feature(buffer, 6);
        features[9] = ccl_get_feature(buffer, 7);
+       if(use_time) {
+               features[10] = pixel.z;
+       }
        if(mean) {
-               for(int i = 0; i < DENOISE_FEATURES; i++)
+               for(int i = 0; i < (use_time? 11 : 10); i++) {
                        features[i] -= mean[i];
+               }
        }
 }
 
-ccl_device_inline void filter_get_feature_scales(int2 pixel,
+ccl_device_inline void filter_get_feature_scales(int3 pixel,
                                                  const ccl_global float *ccl_restrict buffer,
                                                  float *scales,
+                                                 bool use_time,
                                                  const float *ccl_restrict mean,
                                                  int pass_stride)
 {
@@ -66,13 +75,19 @@ ccl_device_inline void filter_get_feature_scales(int2 pixel,
        scales[5] = len_squared(make_float3(ccl_get_feature(buffer, 5) - mean[7],
                                            ccl_get_feature(buffer, 6) - mean[8],
                                            ccl_get_feature(buffer, 7) - mean[9]));
+       if(use_time) {
+               scales[6] = fabsf(pixel.z - mean[10]);
+       }
 }
 
-ccl_device_inline void filter_calculate_scale(float *scale)
+ccl_device_inline void filter_calculate_scale(float *scale, bool use_time)
 {
        scale[0] = 1.0f/max(scale[0], 0.01f);
        scale[1] = 1.0f/max(scale[1], 0.01f);
        scale[2] = 1.0f/max(scale[2], 0.01f);
+       if(use_time) {
+               scale[10] = 1.0f/max(scale[6], 0.01f);
+       }
        scale[6] = 1.0f/max(scale[4], 0.01f);
        scale[7] = scale[8] = scale[9] = 1.0f/max(sqrtf(scale[5]), 0.01f);
        scale[3] = scale[4] = scale[5] = 1.0f/max(sqrtf(scale[3]), 0.01f);
@@ -89,36 +104,46 @@ ccl_device_inline void design_row_add(float *design_row,
                                       const ccl_global float *ccl_restrict transform,
                                       int stride,
                                       int row,
-                                      float feature)
+                                      float feature,
+                                      int transform_row_stride)
 {
        for(int i = 0; i < rank; i++) {
-               design_row[1+i] += transform[(row*DENOISE_FEATURES + i)*stride]*feature;
+               design_row[1+i] += transform[(row*transform_row_stride + i)*stride]*feature;
        }
 }
 
 /* Fill the design row. */
-ccl_device_inline void filter_get_design_row_transform(int2 p_pixel,
+ccl_device_inline void filter_get_design_row_transform(int3 p_pixel,
                                                        const ccl_global float *ccl_restrict p_buffer,
-                                                       int2 q_pixel,
+                                                       int3 q_pixel,
                                                        const ccl_global float *ccl_restrict q_buffer,
                                                        int pass_stride,
                                                        int rank,
                                                        float *design_row,
                                                        const ccl_global float *ccl_restrict transform,
-                                                       int stride)
+                                                       int stride,
+                                                       bool use_time)
 {
+       int num_features = use_time? 11 : 10;
+
        design_row[0] = 1.0f;
        math_vector_zero(design_row+1, rank);
-       design_row_add(design_row, rank, transform, stride, 0, q_pixel.x - p_pixel.x);
-       design_row_add(design_row, rank, transform, stride, 1, q_pixel.y - p_pixel.y);
-       design_row_add(design_row, rank, transform, stride, 2, fabsf(ccl_get_feature(q_buffer, 0)) - fabsf(ccl_get_feature(p_buffer, 0)));
-       design_row_add(design_row, rank, transform, stride, 3, ccl_get_feature(q_buffer, 1) - ccl_get_feature(p_buffer, 1));
-       design_row_add(design_row, rank, transform, stride, 4, ccl_get_feature(q_buffer, 2) - ccl_get_feature(p_buffer, 2));
-       design_row_add(design_row, rank, transform, stride, 5, ccl_get_feature(q_buffer, 3) - ccl_get_feature(p_buffer, 3));
-       design_row_add(design_row, rank, transform, stride, 6, ccl_get_feature(q_buffer, 4) - ccl_get_feature(p_buffer, 4));
-       design_row_add(design_row, rank, transform, stride, 7, ccl_get_feature(q_buffer, 5) - ccl_get_feature(p_buffer, 5));
-       design_row_add(design_row, rank, transform, stride, 8, ccl_get_feature(q_buffer, 6) - ccl_get_feature(p_buffer, 6));
-       design_row_add(design_row, rank, transform, stride, 9, ccl_get_feature(q_buffer, 7) - ccl_get_feature(p_buffer, 7));
+
+#define DESIGN_ROW_ADD(I, F) design_row_add(design_row, rank, transform, stride, I, F, num_features);
+       DESIGN_ROW_ADD(0, q_pixel.x - p_pixel.x);
+       DESIGN_ROW_ADD(1, q_pixel.y - p_pixel.y);
+       DESIGN_ROW_ADD(2, fabsf(ccl_get_feature(q_buffer, 0)) - fabsf(ccl_get_feature(p_buffer, 0)));
+       DESIGN_ROW_ADD(3,       ccl_get_feature(q_buffer, 1)  -       ccl_get_feature(p_buffer, 1));
+       DESIGN_ROW_ADD(4,       ccl_get_feature(q_buffer, 2)  -       ccl_get_feature(p_buffer, 2));
+       DESIGN_ROW_ADD(5,       ccl_get_feature(q_buffer, 3)  -       ccl_get_feature(p_buffer, 3));
+       DESIGN_ROW_ADD(6,       ccl_get_feature(q_buffer, 4)  -       ccl_get_feature(p_buffer, 4));
+       DESIGN_ROW_ADD(7,       ccl_get_feature(q_buffer, 5)  -       ccl_get_feature(p_buffer, 5));
+       DESIGN_ROW_ADD(8,       ccl_get_feature(q_buffer, 6)  -       ccl_get_feature(p_buffer, 6));
+       DESIGN_ROW_ADD(9,       ccl_get_feature(q_buffer, 7)  -       ccl_get_feature(p_buffer, 7));
+       if(use_time) {
+               DESIGN_ROW_ADD(10, q_pixel.z - p_pixel.z)
+       }
+#undef DESIGN_ROW_ADD
 }
 
 CCL_NAMESPACE_END
index 3ddd871..5dd001f 100644 (file)
@@ -20,26 +20,33 @@ CCL_NAMESPACE_BEGIN
 
 /* Loop over the pixels in the range [low.x, high.x) x [low.y, high.y), 4 at a time.
  * pixel_buffer always points to the first of the 4 current pixel in the first pass.
- * x4 and y4 contain the coordinates of the four pixels, active_pixels contains a mask that's set for all pixels within the window. */
+ * x4 and y4 contain the coordinates of the four pixels, active_pixels contains a mask that's set for all pixels within the window.
+ * Repeat the loop for every secondary frame if there are any. */
+#define FOR_PIXEL_WINDOW_SSE     for(int frame = 0; frame < tile_info->num_frames; frame++) { \
+                                     pixel.z = tile_info->frames[frame]; \
+                                     pixel_buffer = buffer + (low.y - rect.y)*buffer_w + (low.x - rect.x) + frame*frame_stride; \
+                                     float4 t4 = make_float4(pixel.z); \
+                                     for(pixel.y = low.y; pixel.y < high.y; pixel.y++) { \
+                                         float4 y4 = make_float4(pixel.y); \
+                                         for(pixel.x = low.x; pixel.x < high.x; pixel.x += 4, pixel_buffer += 4) { \
+                                             float4 x4 = make_float4(pixel.x) + make_float4(0.0f, 1.0f, 2.0f, 3.0f); \
+                                             int4 active_pixels = x4 < make_float4(high.x);
 
-#define FOR_PIXEL_WINDOW_SSE     pixel_buffer = buffer + (low.y - rect.y)*buffer_w + (low.x - rect.x); \
-                                 for(pixel.y = low.y; pixel.y < high.y; pixel.y++) { \
-                                     float4 y4 = make_float4(pixel.y); \
-                                     for(pixel.x = low.x; pixel.x < high.x; pixel.x += 4, pixel_buffer += 4) { \
-                                         float4 x4 = make_float4(pixel.x) + make_float4(0.0f, 1.0f, 2.0f, 3.0f); \
-                                         int4 active_pixels = x4 < make_float4(high.x);
-
-#define END_FOR_PIXEL_WINDOW_SSE     } \
-                                     pixel_buffer += buffer_w - (pixel.x - low.x); \
+#define END_FOR_PIXEL_WINDOW_SSE         } \
+                                         pixel_buffer += buffer_w - (high.x - low.x); \
+                                     } \
                                  }
 
-ccl_device_inline void filter_get_features_sse(float4 x, float4 y,
+ccl_device_inline void filter_get_features_sse(float4 x, float4 y, float4 t,
                                                int4 active_pixels,
                                                const float *ccl_restrict buffer,
                                                float4 *features,
+                                               bool use_time,
                                                const float4 *ccl_restrict mean,
                                                int pass_stride)
 {
+       int num_features = use_time? 11 : 10;
+
        features[0] = x;
        features[1] = y;
        features[2] = fabs(ccl_get_feature_sse(0));
@@ -50,18 +57,25 @@ ccl_device_inline void filter_get_features_sse(float4 x, float4 y,
        features[7] = ccl_get_feature_sse(5);
        features[8] = ccl_get_feature_sse(6);
        features[9] = ccl_get_feature_sse(7);
+       if(use_time) {
+               features[10] = t;
+       }
+
        if(mean) {
-               for(int i = 0; i < DENOISE_FEATURES; i++)
+               for(int i = 0; i < num_features; i++) {
                        features[i] = features[i] - mean[i];
+               }
        }
-       for(int i = 0; i < DENOISE_FEATURES; i++)
+       for(int i = 0; i < num_features; i++) {
                features[i] = mask(active_pixels, features[i]);
+       }
 }
 
-ccl_device_inline void filter_get_feature_scales_sse(float4 x, float4 y,
+ccl_device_inline void filter_get_feature_scales_sse(float4 x, float4 y, float4 t,
                                                      int4 active_pixels,
                                                      const float *ccl_restrict buffer,
                                                      float4 *scales,
+                                                     bool use_time,
                                                      const float4 *ccl_restrict mean,
                                                      int pass_stride)
 {
@@ -75,15 +89,22 @@ ccl_device_inline void filter_get_feature_scales_sse(float4 x, float4 y,
        scales[5] = sqr(ccl_get_feature_sse(5) - mean[7]) +
                    sqr(ccl_get_feature_sse(6) - mean[8]) +
                    sqr(ccl_get_feature_sse(7) - mean[9]);
-       for(int i = 0; i < 6; i++)
+       if(use_time) {
+               scales[6] = fabs(t - mean[10]);
+       }
+
+       for(int i = 0; i < (use_time? 7 : 6); i++)
                scales[i] = mask(active_pixels, scales[i]);
 }
 
-ccl_device_inline void filter_calculate_scale_sse(float4 *scale)
+ccl_device_inline void filter_calculate_scale_sse(float4 *scale, bool use_time)
 {
        scale[0] = rcp(max(reduce_max(scale[0]), make_float4(0.01f)));
        scale[1] = rcp(max(reduce_max(scale[1]), make_float4(0.01f)));
        scale[2] = rcp(max(reduce_max(scale[2]), make_float4(0.01f)));
+       if(use_time) {
+               scale[10] = rcp(max(reduce_max(scale[6]), make_float4(0.01f)));;
+       }
        scale[6] = rcp(max(reduce_max(scale[4]), make_float4(0.01f)));
        scale[7] = scale[8] = scale[9] = rcp(max(reduce_max(sqrt(scale[5])), make_float4(0.01f)));
        scale[3] = scale[4] = scale[5] = rcp(max(reduce_max(sqrt(scale[3])), make_float4(0.01f)));
index af73c0d..9eb3c60 100644 (file)
@@ -22,10 +22,12 @@ CCL_NAMESPACE_BEGIN
 ccl_device_inline void kernel_filter_nlm_calc_difference(int dx, int dy,
                                                          const float *ccl_restrict weight_image,
                                                          const float *ccl_restrict variance_image,
+                                                         const float *ccl_restrict scale_image,
                                                          float *difference_image,
                                                          int4 rect,
                                                          int stride,
                                                          int channel_offset,
+                                                         int frame_offset,
                                                          float a,
                                                          float k_2)
 {
@@ -38,16 +40,24 @@ ccl_device_inline void kernel_filter_nlm_calc_difference(int dx, int dy,
 
        for(int y = rect.y; y < rect.w; y++) {
                int idx_p = y*stride + aligned_lowx;
-               int idx_q = (y+dy)*stride + aligned_lowx + dx;
+               int idx_q = (y+dy)*stride + aligned_lowx + dx + frame_offset;
                for(int x = aligned_lowx; x < rect.z; x += 4, idx_p += 4, idx_q += 4) {
                        float4 diff = make_float4(0.0f);
+                       float4 scale_fac;
+                       if(scale_image) {
+                               scale_fac = clamp(load4_a(scale_image, idx_p) / load4_u(scale_image, idx_q),
+                                                 make_float4(0.25f), make_float4(4.0f));
+                       }
+                       else {
+                               scale_fac = make_float4(1.0f);
+                       }
                        for(int c = 0, chan_ofs = 0; c < numChannels; c++, chan_ofs += channel_offset) {
                                /* idx_p is guaranteed to be aligned, but idx_q isn't. */
                                float4 color_p = load4_a(weight_image, idx_p + chan_ofs);
-                               float4 color_q = load4_u(weight_image, idx_q + chan_ofs);
+                               float4 color_q = scale_fac*load4_u(weight_image, idx_q + chan_ofs);
                                float4 cdiff = color_p - color_q;
                                float4 var_p = load4_a(variance_image, idx_p + chan_ofs);
-                               float4 var_q = load4_u(variance_image, idx_q + chan_ofs);
+                               float4 var_q = sqr(scale_fac)*load4_u(variance_image, idx_q + chan_ofs);
                                diff += (cdiff*cdiff - a*(var_p + min(var_p, var_q))) / (make_float4(1e-8f) + k_2*(var_p+var_q));
                        }
                        load4_a(difference_image, idx_p) = diff*channel_fac;
@@ -143,6 +153,7 @@ ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy,
                                                        float *out_image,
                                                        float *accum_image,
                                                        int4 rect,
+                                                       int channel_offset,
                                                        int stride,
                                                        int f)
 {
@@ -160,13 +171,18 @@ ccl_device_inline void kernel_filter_nlm_update_output(int dx, int dy,
                        load4_a(accum_image, idx_p) += mask(active, weight);
 
                        float4 val = load4_u(image, idx_q);
+                       if(channel_offset) {
+                               val += load4_u(image, idx_q + channel_offset);
+                               val += load4_u(image, idx_q + 2*channel_offset);
+                               val *= 1.0f/3.0f;
+                       }
 
                        load4_a(out_image, idx_p) += mask(active, weight*val);
                }
        }
 }
 
-ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy,
+ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy, int t,
                                                            const float *ccl_restrict difference_image,
                                                            const float *ccl_restrict buffer,
                                                            float *transform,
@@ -176,7 +192,9 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy,
                                                            int4 rect,
                                                            int4 filter_window,
                                                            int stride, int f,
-                                                           int pass_stride)
+                                                           int pass_stride,
+                                                           int frame_offset,
+                                                           bool use_time)
 {
        int4 clip_area = rect_clip(rect, filter_window);
        /* fy and fy are in filter-window-relative coordinates, while x and y are in feature-window-relative coordinates. */
@@ -197,9 +215,11 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int dx, int dy,
                        int    *l_rank = rank + storage_ofs;
 
                        kernel_filter_construct_gramian(x, y, 1,
-                                                       dx, dy,
+                                                       dx, dy, t,
                                                        stride,
                                                        pass_stride,
+                                                       frame_offset,
+                                                       use_time,
                                                        buffer,
                                                        l_transform, l_rank,
                                                        weight, l_XtWX, l_XtWY, 0);
index 058afb3..1263639 100644 (file)
@@ -78,17 +78,26 @@ ccl_device_inline void kernel_filter_nlm_calc_difference(int x, int y,
                                                          int dx, int dy,
                                                          const ccl_global float *ccl_restrict weight_image,
                                                          const ccl_global float *ccl_restrict variance_image,
+                                                         const ccl_global float *ccl_restrict scale_image,
                                                          ccl_global float *difference_image,
                                                          int4 rect, int stride,
                                                          int channel_offset,
+                                                         int frame_offset,
                                                          float a, float k_2)
 {
-       float diff = 0.0f;
+       int idx_p = y*stride + x, idx_q = (y+dy)*stride + (x+dx) + frame_offset;
        int numChannels = channel_offset? 3 : 1;
-       for(int c = 0; c < numChannels; c++) {
-               float cdiff = weight_image[c*channel_offset + y*stride + x] - weight_image[c*channel_offset + (y+dy)*stride + (x+dx)];
-               float pvar = variance_image[c*channel_offset + y*stride + x];
-               float qvar = variance_image[c*channel_offset + (y+dy)*stride + (x+dx)];
+
+       float diff = 0.0f;
+       float scale_fac = 1.0f;
+       if(scale_image) {
+               scale_fac = clamp(scale_image[idx_p] / scale_image[idx_q], 0.25f, 4.0f);
+       }
+
+       for(int c = 0; c < numChannels; c++, idx_p += channel_offset, idx_q += channel_offset) {
+               float cdiff = weight_image[idx_p] - scale_fac*weight_image[idx_q];
+               float pvar = variance_image[idx_p];
+               float qvar = sqr(scale_fac)*variance_image[idx_q];
                diff += (cdiff*cdiff - a*(pvar + min(pvar, qvar))) / (1e-8f + k_2*(pvar+qvar));
        }
        if(numChannels > 1) {
@@ -133,7 +142,8 @@ ccl_device_inline void kernel_filter_nlm_update_output(int x, int y,
                                                        const ccl_global float *ccl_restrict image,
                                                        ccl_global float *out_image,
                                                        ccl_global float *accum_image,
-                                                       int4 rect, int stride, int f)
+                                                       int4 rect, int channel_offset,
+                                                       int stride, int f)
 {
        float sum = 0.0f;
        const int low = max(rect.x, x-f);
@@ -142,17 +152,26 @@ ccl_device_inline void kernel_filter_nlm_update_output(int x, int y,
                sum += difference_image[y*stride + x1];
        }
        sum *= 1.0f/(high-low);
+
+       int idx_p = y*stride + x, idx_q = (y+dy)*stride + (x+dx);
        if(out_image) {
-               atomic_add_and_fetch_float(accum_image + y*stride + x, sum);
-               atomic_add_and_fetch_float(out_image + y*stride + x, sum*image[(y+dy)*stride + (x+dx)]);
+               atomic_add_and_fetch_float(accum_image + idx_p, sum);
+
+               float val = image[idx_q];
+               if(channel_offset) {
+                       val += image[idx_q + channel_offset];
+                       val += image[idx_q + 2*channel_offset];
+                       val *= 1.0f/3.0f;
+               }
+               atomic_add_and_fetch_float(out_image + idx_p, sum*val);
        }
        else {
-               accum_image[y*stride + x] = sum;
+               accum_image[idx_p] = sum;
        }
 }
 
 ccl_device_inline void kernel_filter_nlm_construct_gramian(int x, int y,
-                                                           int dx, int dy,
+                                                           int dx, int dy, int t,
                                                            const ccl_global float *ccl_restrict difference_image,
                                                            const ccl_global float *ccl_restrict buffer,
                                                            const ccl_global float *ccl_restrict transform,
@@ -163,6 +182,8 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int x, int y,
                                                            int4 filter_window,
                                                            int stride, int f,
                                                            int pass_stride,
+                                                           int frame_offset,
+                                                           bool use_time,
                                                            int localIdx)
 {
        const int low = max(rect.x, x-f);
@@ -183,9 +204,11 @@ ccl_device_inline void kernel_filter_nlm_construct_gramian(int x, int y,
 
        kernel_filter_construct_gramian(x, y,
                                        rect_size(filter_window),
-                                       dx, dy,
+                                       dx, dy, t,
                                        stride,
                                        pass_stride,
+                                       frame_offset,
+                                       use_time,
                                        buffer,
                                        transform, rank,
                                        weight, XtWX, XtWY,
index 3507f80..e24f4fe 100644 (file)
@@ -84,6 +84,7 @@ ccl_device void kernel_filter_get_feature(int sample,
                                           int x, int y,
                                           ccl_global float *mean,
                                           ccl_global float *variance,
+                                          float scale,
                                           int4 rect, int buffer_pass_stride,
                                           int buffer_denoising_offset)
 {
@@ -95,18 +96,38 @@ ccl_device void kernel_filter_get_feature(int sample,
        int buffer_w = align_up(rect.z - rect.x, 4);
        int idx = (y-rect.y)*buffer_w + (x - rect.x);
 
-       mean[idx] = center_buffer[m_offset] / sample;
-       if(sample > 1) {
-               /* Approximate variance as E[x^2] - 1/N * (E[x])^2, since online variance
-                * update does not work efficiently with atomics in the kernel. */
-               variance[idx] = max(0.0f, (center_buffer[v_offset] - mean[idx]*mean[idx]*sample) / (sample * (sample-1)));
-       }
-       else {
-               /* Can't compute variance with single sample, just set it very high. */
-               variance[idx] = 1e10f;
+       float val = scale * center_buffer[m_offset];
+       mean[idx] = val;
+
+       if(v_offset >= 0) {
+               if(sample > 1) {
+                       /* Approximate variance as E[x^2] - 1/N * (E[x])^2, since online variance
+                        * update does not work efficiently with atomics in the kernel. */
+                       variance[idx] = max(0.0f, (center_buffer[v_offset] - val*val*sample) / (sample * (sample-1)));
+               }
+               else {
+                       /* Can't compute variance with single sample, just set it very high. */
+                       variance[idx] = 1e10f;
+               }
        }
 }
 
+ccl_device void kernel_filter_write_feature(int sample,
+                                            int x, int y,
+                                            int4 buffer_params,
+                                            ccl_global float *from,
+                                            ccl_global float *buffer,
+                                            int out_offset,
+                                            int4 rect)
+{
+       ccl_global float *combined_buffer = buffer + (y*buffer_params.y + x + buffer_params.x)*buffer_params.z;
+
+       int buffer_w = align_up(rect.z - rect.x, 4);
+       int idx = (y-rect.y)*buffer_w + (x - rect.x);
+
+       combined_buffer[out_offset] = from[idx];
+}
+
 ccl_device void kernel_filter_detect_outliers(int x, int y,
                                               ccl_global float *image,
                                               ccl_global float *variance,
@@ -119,6 +140,7 @@ ccl_device void kernel_filter_detect_outliers(int x, int y,
 
        int n = 0;
        float values[25];
+       float pixel_variance, max_variance = 0.0f;
        for(int y1 = max(y-2, rect.y); y1 < min(y+3, rect.w); y1++) {
                for(int x1 = max(x-2, rect.x); x1 < min(x+3, rect.z); x1++) {
                        int idx = (y1-rect.y)*buffer_w + (x1-rect.x);
@@ -138,15 +160,31 @@ ccl_device void kernel_filter_detect_outliers(int x, int y,
                        /* Insert L. */
                        values[i] = L;
                        n++;
+
+                       float3 pixel_var = make_float3(variance[idx], variance[idx+pass_stride], variance[idx+2*pass_stride]);
+                       float var = average(pixel_var);
+                       if((x1 == x) && (y1 == y)) {
+                               pixel_variance = (pixel_var.x < 0.0f || pixel_var.y < 0.0f || pixel_var.z < 0.0f)? -1.0f : var;
+                       }
+                       else {
+                               max_variance = max(max_variance, var);
+                       }
                }
        }
 
+       max_variance += 1e-4f;
+
        int idx = (y-rect.y)*buffer_w + (x-rect.x);
        float3 color = make_float3(image[idx], image[idx+pass_stride], image[idx+2*pass_stride]);
        color = max(color, make_float3(0.0f, 0.0f, 0.0f));
        float L = average(color);
 
        float ref = 2.0f*values[(int)(n*0.75f)];
+
+       /* Slightly offset values to avoid false positives in (almost) black areas. */
+       max_variance += 1e-5f;
+       ref -= 1e-5f;
+
        if(L > ref) {
                /* The pixel appears to be an outlier.
                 * However, it may just be a legitimate highlight. Therefore, it is checked how likely it is that the pixel
@@ -154,16 +192,24 @@ ccl_device void kernel_filter_detect_outliers(int x, int y,
                 * If the reference is within the 3-sigma interval, the pixel is assumed to be a statistical outlier.
                 * Otherwise, it is very unlikely that the pixel should be darker, which indicates a legitimate highlight.
                 */
-               float stddev = sqrtf(average(make_float3(variance[idx], variance[idx+pass_stride], variance[idx+2*pass_stride])));
-               if(L - 3*stddev < ref) {
-                       /* The pixel is an outlier, so negate the depth value to mark it as one.
-                        * Also, scale its brightness down to the outlier threshold to avoid trouble with the NLM weights. */
+
+               if(pixel_variance < 0.0f || pixel_variance > 9.0f * max_variance) {
                        depth[idx] = -depth[idx];
-                       float fac = ref/L;
-                       color *= fac;
-                       variance[idx              ] *= fac*fac;
-                       variance[idx + pass_stride] *= fac*fac;
-                       variance[idx+2*pass_stride] *= fac*fac;
+                       color *= ref/L;
+                       variance[idx] = variance[idx + pass_stride] = variance[idx + 2*pass_stride] = max_variance;
+               }
+               else {
+                       float stddev = sqrtf(pixel_variance);
+                       if(L - 3*stddev < ref) {
+                               /* The pixel is an outlier, so negate the depth value to mark it as one.
+                               * Also, scale its brightness down to the outlier threshold to avoid trouble with the NLM weights. */
+                               depth[idx] = -depth[idx];
+                               float fac = ref/L;
+                               color *= fac;
+                               variance[idx              ] *= fac*fac;
+                               variance[idx + pass_stride] *= fac*fac;
+                               variance[idx+2*pass_stride] *= fac*fac;
+                       }
                }
        }
        out[idx              ] = color.x;
index 58740d5..31a7487 100644 (file)
@@ -18,9 +18,11 @@ CCL_NAMESPACE_BEGIN
 
 ccl_device_inline void kernel_filter_construct_gramian(int x, int y,
                                                        int storage_stride,
-                                                       int dx, int dy,
+                                                       int dx, int dy, int t,
                                                        int buffer_stride,
                                                        int pass_stride,
+                                                       int frame_offset,
+                                                       bool use_time,
                                                        const ccl_global float *ccl_restrict buffer,
                                                        const ccl_global float *ccl_restrict transform,
                                                        ccl_global int *rank,
@@ -34,7 +36,7 @@ ccl_device_inline void kernel_filter_construct_gramian(int x, int y,
        }
 
        int p_offset =  y     * buffer_stride +  x;
-       int q_offset = (y+dy) * buffer_stride + (x+dx);
+       int q_offset = (y+dy) * buffer_stride + (x+dx) + frame_offset;
 
 #ifdef __KERNEL_GPU__
        const int stride = storage_stride;
@@ -57,9 +59,9 @@ ccl_device_inline void kernel_filter_construct_gramian(int x, int y,
                return;
        }
 
-       filter_get_design_row_transform(make_int2(x, y),       buffer + p_offset,
-                                       make_int2(x+dx, y+dy), buffer + q_offset,
-                                       pass_stride, *rank, design_row, transform, stride);
+       filter_get_design_row_transform(make_int3(x, y, t),       buffer + p_offset,
+                                       make_int3(x+dx, y+dy, t), buffer + q_offset,
+                                       pass_stride, *rank, design_row, transform, stride, use_time);
 
 #ifdef __KERNEL_GPU__
        math_trimatrix_add_gramian_strided(XtWX, (*rank)+1, design_row, weight, stride);
@@ -108,11 +110,13 @@ ccl_device_inline void kernel_filter_finalize(int x, int y,
        final_color = max(final_color, make_float3(0.0f, 0.0f, 0.0f));
 
        ccl_global float *combined_buffer = buffer + (y*buffer_params.y + x + buffer_params.x)*buffer_params.z;
-       final_color *= sample;
-       if(buffer_params.w) {
-               final_color.x += combined_buffer[buffer_params.w+0];
-               final_color.y += combined_buffer[buffer_params.w+1];
-               final_color.z += combined_buffer[buffer_params.w+2];
+       if(buffer_params.w >= 0) {
+               final_color *= sample;
+               if(buffer_params.w > 0) {
+                       final_color.x += combined_buffer[buffer_params.w+0];
+                       final_color.y += combined_buffer[buffer_params.w+1];
+                       final_color.z += combined_buffer[buffer_params.w+2];
+               }
        }
        combined_buffer[0] = final_color.x;
        combined_buffer[1] = final_color.y;
index a5f87c0..94e27bb 100644 (file)
 CCL_NAMESPACE_BEGIN
 
 ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buffer,
+                                                  CCL_FILTER_TILE_INFO,
                                                   int x, int y, int4 rect,
-                                                  int pass_stride,
+                                                  int pass_stride, int frame_stride,
+                                                  bool use_time,
                                                   float *transform, int *rank,
                                                   int radius, float pca_threshold)
 {
@@ -26,59 +28,58 @@ ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buff
 
        float features[DENOISE_FEATURES];
 
-       /* Temporary storage, used in different steps of the algorithm. */
-       float tempmatrix[DENOISE_FEATURES*DENOISE_FEATURES];
-       float tempvector[2*DENOISE_FEATURES];
        const float *ccl_restrict pixel_buffer;
-       int2 pixel;
+       int3 pixel;
+
+       int num_features = use_time? 11 : 10;
 
        /* === Calculate denoising window. === */
        int2 low  = make_int2(max(rect.x, x - radius),
                              max(rect.y, y - radius));
        int2 high = make_int2(min(rect.z, x + radius + 1),
                              min(rect.w, y + radius + 1));
-       int num_pixels = (high.y - low.y) * (high.x - low.x);
+       int num_pixels = (high.y - low.y) * (high.x - low.x) * tile_info->num_frames;
 
        /* === Shift feature passes to have mean 0. === */
        float feature_means[DENOISE_FEATURES];
-       math_vector_zero(feature_means, DENOISE_FEATURES);
+       math_vector_zero(feature_means, num_features);
        FOR_PIXEL_WINDOW {
-               filter_get_features(pixel, pixel_buffer, features, NULL, pass_stride);
-               math_vector_add(feature_means, features, DENOISE_FEATURES);
+               filter_get_features(pixel, pixel_buffer, features, use_time, NULL, pass_stride);
+               math_vector_add(feature_means, features, num_features);
        } END_FOR_PIXEL_WINDOW
 
-       math_vector_scale(feature_means, 1.0f / num_pixels, DENOISE_FEATURES);
+       math_vector_scale(feature_means, 1.0f / num_pixels, num_features);
 
        /* === Scale the shifted feature passes to a range of [-1; 1], will be baked into the transform later. === */
-       float *feature_scale = tempvector;
-       math_vector_zero(feature_scale, DENOISE_FEATURES);
+       float feature_scale[DENOISE_FEATURES];
+       math_vector_zero(feature_scale, num_features);
 
        FOR_PIXEL_WINDOW {
-               filter_get_feature_scales(pixel, pixel_buffer, features, feature_means, pass_stride);
-               math_vector_max(feature_scale, features, DENOISE_FEATURES);
+               filter_get_feature_scales(pixel, pixel_buffer, features, use_time, feature_means, pass_stride);
+               math_vector_max(feature_scale, features, num_features);
        } END_FOR_PIXEL_WINDOW
 
-       filter_calculate_scale(feature_scale);
+       filter_calculate_scale(feature_scale, use_time);
 
        /* === Generate the feature transformation. ===
-        * This transformation maps the DENOISE_FEATURES-dimentional feature space to a reduced feature (r-feature) space
+        * This transformation maps the num_features-dimentional feature space to a reduced feature (r-feature) space
         * which generally has fewer dimensions. This mainly helps to prevent overfitting. */
-       float* feature_matrix = tempmatrix;
-       math_matrix_zero(feature_matrix, DENOISE_FEATURES);
+       float feature_matrix[DENOISE_FEATURES*DENOISE_FEATURES];
+       math_matrix_zero(feature_matrix, num_features);
        FOR_PIXEL_WINDOW {
-               filter_get_features(pixel, pixel_buffer, features, feature_means, pass_stride);
-               math_vector_mul(features, feature_scale, DENOISE_FEATURES);
-               math_matrix_add_gramian(feature_matrix, DENOISE_FEATURES, features, 1.0f);
+               filter_get_features(pixel, pixel_buffer, features, use_time, feature_means, pass_stride);
+               math_vector_mul(features, feature_scale, num_features);
+               math_matrix_add_gramian(feature_matrix, num_features, features, 1.0f);
        } END_FOR_PIXEL_WINDOW
 
-       math_matrix_jacobi_eigendecomposition(feature_matrix, transform, DENOISE_FEATURES, 1);
+       math_matrix_jacobi_eigendecomposition(feature_matrix, transform, num_features, 1);
        *rank = 0;
        /* Prevent overfitting when a small window is used. */
-       int max_rank = min(DENOISE_FEATURES, num_pixels/3);
+       int max_rank = min(num_features, num_pixels/3);
        if(pca_threshold < 0.0f) {
                float threshold_energy = 0.0f;
-               for(int i = 0; i < DENOISE_FEATURES; i++) {
-                       threshold_energy += feature_matrix[i*DENOISE_FEATURES+i];
+               for(int i = 0; i < num_features; i++) {
+                       threshold_energy += feature_matrix[i*num_features+i];
                }
                threshold_energy *= 1.0f - (-pca_threshold);
 
@@ -86,13 +87,13 @@ ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buff
                for(int i = 0; i < max_rank; i++, (*rank)++) {
                        if(i >= 2 && reduced_energy >= threshold_energy)
                                break;
-                       float s = feature_matrix[i*DENOISE_FEATURES+i];
+                       float s = feature_matrix[i*num_features+i];
                        reduced_energy += s;
                }
        }
        else {
                for(int i = 0; i < max_rank; i++, (*rank)++) {
-                       float s = feature_matrix[i*DENOISE_FEATURES+i];
+                       float s = feature_matrix[i*num_features+i];
                        if(i >= 2 && sqrtf(s) < pca_threshold)
                                break;
                }
@@ -100,9 +101,9 @@ ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buff
 
        /* Bake the feature scaling into the transformation matrix. */
        for(int i = 0; i < (*rank); i++) {
-               math_vector_mul(transform + i*DENOISE_FEATURES, feature_scale, DENOISE_FEATURES);
+               math_vector_mul(transform + i*num_features, feature_scale, num_features);
        }
-       math_matrix_transpose(transform, DENOISE_FEATURES, 1);
+       math_matrix_transpose(transform, num_features, 1);
 }
 
 CCL_NAMESPACE_END
index 83a1222..ed8ddcb 100644 (file)
 CCL_NAMESPACE_BEGIN
 
 ccl_device void kernel_filter_construct_transform(const ccl_global float *ccl_restrict buffer,
+                                                  CCL_FILTER_TILE_INFO,
                                                   int x, int y, int4 rect,
-                                                  int pass_stride,
+                                                  int pass_stride, int frame_stride,
+                                                  bool use_time,
                                                   ccl_global float *transform,
                                                   ccl_global int *rank,
                                                   int radius, float pca_threshold,
@@ -33,60 +35,62 @@ ccl_device void kernel_filter_construct_transform(const ccl_global float *ccl_re
        float features[DENOISE_FEATURES];
 #endif
 
+       int num_features = use_time? 11 : 10;
+
        /* === Calculate denoising window. === */
        int2 low  = make_int2(max(rect.x, x - radius),
                              max(rect.y, y - radius));
        int2 high = make_int2(min(rect.z, x + radius + 1),
                              min(rect.w, y + radius + 1));
-       int num_pixels = (high.y - low.y) * (high.x - low.x);
+       int num_pixels = (high.y - low.y) * (high.x - low.x) * tile_info->num_frames;
        const ccl_global float *ccl_restrict pixel_buffer;
-       int2 pixel;
+       int3 pixel;
 
 
 
 
        /* === Shift feature passes to have mean 0. === */
        float feature_means[DENOISE_FEATURES];
-       math_vector_zero(feature_means, DENOISE_FEATURES);
+       math_vector_zero(feature_means, num_features);
        FOR_PIXEL_WINDOW {
-               filter_get_features(pixel, pixel_buffer, features, NULL, pass_stride);
-               math_vector_add(feature_means, features, DENOISE_FEATURES);
+               filter_get_features(pixel, pixel_buffer, features, use_time, NULL, pass_stride);
+               math_vector_add(feature_means, features, num_features);
        } END_FOR_PIXEL_WINDOW
 
-       math_vector_scale(feature_means, 1.0f / num_pixels, DENOISE_FEATURES);
+       math_vector_scale(feature_means, 1.0f / num_pixels, num_features);
 
        /* === Scale the shifted feature passes to a range of [-1; 1], will be baked into the transform later. === */
        float feature_scale[DENOISE_FEATURES];
-       math_vector_zero(feature_scale, DENOISE_FEATURES);
+       math_vector_zero(feature_scale, num_features);
 
        FOR_PIXEL_WINDOW {
-               filter_get_feature_scales(pixel, pixel_buffer, features, feature_means, pass_stride);
-               math_vector_max(feature_scale, features, DENOISE_FEATURES);
+               filter_get_feature_scales(pixel, pixel_buffer, features, use_time, feature_means, pass_stride);
+               math_vector_max(feature_scale, features, num_features);
        } END_FOR_PIXEL_WINDOW
 
-       filter_calculate_scale(feature_scale);
+       filter_calculate_scale(feature_scale, use_time);
 
 
 
        /* === Generate the feature transformation. ===
-        * This transformation maps the DENOISE_FEATURES-dimentional feature space to a reduced feature (r-feature) space
+        * This transformation maps the num_features-dimentional feature space to a reduced feature (r-feature) space
         * which generally has fewer dimensions. This mainly helps to prevent overfitting. */
        float feature_matrix[DENOISE_FEATURES*DENOISE_FEATURES];
-       math_matrix_zero(feature_matrix, DENOISE_FEATURES);
+       math_matrix_zero(feature_matrix, num_features);
        FOR_PIXEL_WINDOW {
-               filter_get_features(pixel, pixel_buffer, features, feature_means, pass_stride);
-               math_vector_mul(features, feature_scale, DENOISE_FEATURES);
-               math_matrix_add_gramian(feature_matrix, DENOISE_FEATURES, features, 1.0f);
+               filter_get_features(pixel, pixel_buffer, features, use_time, feature_means, pass_stride);
+               math_vector_mul(features, feature_scale, num_features);
+               math_matrix_add_gramian(feature_matrix, num_features, features, 1.0f);
        } END_FOR_PIXEL_WINDOW
 
-       math_matrix_jacobi_eigendecomposition(feature_matrix, transform, DENOISE_FEATURES, transform_stride);
+       math_matrix_jacobi_eigendecomposition(feature_matrix, transform, num_features, transform_stride);
        *rank = 0;
        /* Prevent overfitting when a small window is used. */
-       int max_rank = min(DENOISE_FEATURES, num_pixels/3);
+       int max_rank = min(num_features, num_pixels/3);
        if(pca_threshold < 0.0f) {
                float threshold_energy = 0.0f;
-               for(int i = 0; i < DENOISE_FEATURES; i++) {
-                       threshold_energy += feature_matrix[i*DENOISE_FEATURES+i];
+               for(int i = 0; i < num_features; i++) {
+                       threshold_energy += feature_matrix[i*num_features+i];
                }
                threshold_energy *= 1.0f - (-pca_threshold);
 
@@ -94,24 +98,24 @@ ccl_device void kernel_filter_construct_transform(const ccl_global float *ccl_re
                for(int i = 0; i < max_rank; i++, (*rank)++) {
                        if(i >= 2 && reduced_energy >= threshold_energy)
                                break;
-                       float s = feature_matrix[i*DENOISE_FEATURES+i];
+                       float s = feature_matrix[i*num_features+i];
                        reduced_energy += s;
                }
        }
        else {
                for(int i = 0; i < max_rank; i++, (*rank)++) {
-                       float s = feature_matrix[i*DENOISE_FEATURES+i];
+                       float s = feature_matrix[i*num_features+i];
                        if(i >= 2 && sqrtf(s) < pca_threshold)
                                break;
                }
        }
 
-       math_matrix_transpose(transform, DENOISE_FEATURES, transform_stride);
+       math_matrix_transpose(transform, num_features, transform_stride);
 
        /* Bake the feature scaling into the transformation matrix. */
-       for(int i = 0; i < DENOISE_FEATURES; i++) {
+       for(int i = 0; i < num_features; i++) {
                for(int j = 0; j < (*rank); j++) {
-                       transform[(i*DENOISE_FEATURES + j)*transform_stride] *= feature_scale[i];
+                       transform[(i*num_features + j)*transform_stride] *= feature_scale[i];
                }
        }
 }
index 9e65f61..10bd3e4 100644 (file)
 CCL_NAMESPACE_BEGIN
 
 ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buffer,
+                                                  CCL_FILTER_TILE_INFO,
                                                   int x, int y, int4 rect,
-                                                  int pass_stride,
+                                                  int pass_stride, int frame_stride,
+                                                  bool use_time,
                                                   float *transform, int *rank,
                                                   int radius, float pca_threshold)
 {
@@ -26,55 +28,63 @@ ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buff
 
        float4 features[DENOISE_FEATURES];
        const float *ccl_restrict pixel_buffer;
-       int2 pixel;
+       int3 pixel;
 
+       int num_features = use_time? 11 : 10;
+
+       /* === Calculate denoising window. === */
        int2 low  = make_int2(max(rect.x, x - radius),
                              max(rect.y, y - radius));
        int2 high = make_int2(min(rect.z, x + radius + 1),
                              min(rect.w, y + radius + 1));
-       int num_pixels = (high.y - low.y) * (high.x - low.x);
+       int num_pixels = (high.y - low.y) * (high.x - low.x) * tile_info->num_frames;
 
+       /* === Shift feature passes to have mean 0. === */
        float4 feature_means[DENOISE_FEATURES];
-       math_vector_zero_sse(feature_means, DENOISE_FEATURES);
+       math_vector_zero_sse(feature_means, num_features);
        FOR_PIXEL_WINDOW_SSE {
-               filter_get_features_sse(x4, y4, active_pixels, pixel_buffer, features, NULL, pass_stride);
-               math_vector_add_sse(feature_means, DENOISE_FEATURES, features);
+               filter_get_features_sse(x4, y4, t4, active_pixels, pixel_buffer, features, use_time, NULL, pass_stride);
+               math_vector_add_sse(feature_means, num_features, features);
        } END_FOR_PIXEL_WINDOW_SSE
 
        float4 pixel_scale = make_float4(1.0f / num_pixels);
-       for(int i = 0; i < DENOISE_FEATURES; i++) {
+       for(int i = 0; i < num_features; i++) {
                feature_means[i] = reduce_add(feature_means[i]) * pixel_scale;
        }
 
+       /* === Scale the shifted feature passes to a range of [-1; 1], will be baked into the transform later. === */
        float4 feature_scale[DENOISE_FEATURES];
-       math_vector_zero_sse(feature_scale, DENOISE_FEATURES);
+       math_vector_zero_sse(feature_scale, num_features);
        FOR_PIXEL_WINDOW_SSE {
-               filter_get_feature_scales_sse(x4, y4, active_pixels, pixel_buffer, features, feature_means, pass_stride);
-               math_vector_max_sse(feature_scale, features, DENOISE_FEATURES);
+               filter_get_feature_scales_sse(x4, y4, t4, active_pixels, pixel_buffer, features, use_time, feature_means, pass_stride);
+               math_vector_max_sse(feature_scale, features, num_features);
        } END_FOR_PIXEL_WINDOW_SSE
 
-       filter_calculate_scale_sse(feature_scale);
+       filter_calculate_scale_sse(feature_scale, use_time);
 
+       /* === Generate the feature transformation. ===
+        * This transformation maps the num_features-dimentional feature space to a reduced feature (r-feature) space
+        * which generally has fewer dimensions. This mainly helps to prevent overfitting. */
        float4 feature_matrix_sse[DENOISE_FEATURES*DENOISE_FEATURES];
-       math_matrix_zero_sse(feature_matrix_sse, DENOISE_FEATURES);
+       math_matrix_zero_sse(feature_matrix_sse, num_features);
        FOR_PIXEL_WINDOW_SSE {
-               filter_get_features_sse(x4, y4, active_pixels, pixel_buffer, features, feature_means, pass_stride);
-               math_vector_mul_sse(features, DENOISE_FEATURES, feature_scale);
-               math_matrix_add_gramian_sse(feature_matrix_sse, DENOISE_FEATURES, features, make_float4(1.0f));
+               filter_get_features_sse(x4, y4, t4, active_pixels, pixel_buffer, features, use_time, feature_means, pass_stride);
+               math_vector_mul_sse(features, num_features, feature_scale);
+               math_matrix_add_gramian_sse(feature_matrix_sse, num_features, features, make_float4(1.0f));
        } END_FOR_PIXEL_WINDOW_SSE
 
        float feature_matrix[DENOISE_FEATURES*DENOISE_FEATURES];
-       math_matrix_hsum(feature_matrix, DENOISE_FEATURES, feature_matrix_sse);
+       math_matrix_hsum(feature_matrix, num_features, feature_matrix_sse);
 
-       math_matrix_jacobi_eigendecomposition(feature_matrix, transform, DENOISE_FEATURES, 1);
+       math_matrix_jacobi_eigendecomposition(feature_matrix, transform, num_features, 1);
 
        *rank = 0;
        /* Prevent overfitting when a small window is used. */
-       int max_rank = min(DENOISE_FEATURES, num_pixels/3);
+       int max_rank = min(num_features, num_pixels/3);
        if(pca_threshold < 0.0f) {
                float threshold_energy = 0.0f;
-               for(int i = 0; i < DENOISE_FEATURES; i++) {
-                       threshold_energy += feature_matrix[i*DENOISE_FEATURES+i];
+               for(int i = 0; i < num_features; i++) {
+                       threshold_energy += feature_matrix[i*num_features+i];
                }
                threshold_energy *= 1.0f - (-pca_threshold);
 
@@ -82,23 +92,23 @@ ccl_device void kernel_filter_construct_transform(const float *ccl_restrict buff
                for(int i = 0; i < max_rank; i++, (*rank)++) {
                        if(i >= 2 && reduced_energy >= threshold_energy)
                                break;
-                       float s = feature_matrix[i*DENOISE_FEATURES+i];
+                       float s = feature_matrix[i*num_features+i];
                        reduced_energy += s;
                }
        }
        else {
                for(int i = 0; i < max_rank; i++, (*rank)++) {
-                       float s = feature_matrix[i*DENOISE_FEATURES+i];
+                       float s = feature_matrix[i*num_features+i];
                        if(i >= 2 && sqrtf(s) < pca_threshold)
                                break;
                }
        }
 
-       math_matrix_transpose(transform, DENOISE_FEATURES, 1);
+       math_matrix_transpose(transform, num_features, 1);
 
        /* Bake the feature scaling into the transformation matrix. */
-       for(int i = 0; i < DENOISE_FEATURES; i++) {
-               math_vector_scale(transform + i*DENOISE_FEATURES, feature_scale[i][0], *rank);
+       for(int i = 0; i < num_features; i++) {
+               math_vector_scale(transform + i*num_features, feature_scale[i][0], *rank);
        }
 }
 
index 864aa7c..caa0057 100644 (file)
@@ -472,8 +472,17 @@ typedef enum DenoisingPassOffsets {
        DENOISING_PASS_COLOR_VAR          = 23,
        DENOISING_PASS_CLEAN              = 26,
 
+       DENOISING_PASS_PREFILTERED_DEPTH     = 0,
+       DENOISING_PASS_PREFILTERED_NORMAL    = 1,
+       DENOISING_PASS_PREFILTERED_SHADOWING = 4,
+       DENOISING_PASS_PREFILTERED_ALBEDO    = 5,
+       DENOISING_PASS_PREFILTERED_COLOR     = 8,
+       DENOISING_PASS_PREFILTERED_VARIANCE  = 11,
+       DENOISING_PASS_PREFILTERED_INTENSITY = 14,
+
        DENOISING_PASS_SIZE_BASE          = 26,
        DENOISING_PASS_SIZE_CLEAN         = 3,
+       DENOISING_PASS_SIZE_PREFILTERED   = 15,
 } DenoisingPassOffsets;
 
 typedef enum eBakePassFilter {
index e036b53..02c8556 100644 (file)
@@ -37,10 +37,20 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample,
                                                    int y,
                                                    float *mean,
                                                    float *variance,
+                                                   float scale,
                                                    int* prefilter_rect,
                                                    int buffer_pass_stride,
                                                    int buffer_denoising_offset);
 
+void KERNEL_FUNCTION_FULL_NAME(filter_write_feature)(int sample,
+                                                     int x,
+                                                     int y,
+                                                     int *buffer_params,
+                                                     float *from,
+                                                     float *buffer,
+                                                     int out_offset,
+                                                     int* prefilter_rect);
+
 void KERNEL_FUNCTION_FULL_NAME(filter_detect_outliers)(int x, int y,
                                                        ccl_global float *image,
                                                        ccl_global float *variance,
@@ -58,6 +68,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_combine_halves)(int x, int y,
                                                       int r);
 
 void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer,
+                                                           TileInfo *tiles,
                                                            int x,
                                                            int y,
                                                            int storage_ofs,
@@ -65,17 +76,21 @@ void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer,
                                                            int *rank,
                                                            int* rect,
                                                            int pass_stride,
+                                                           int frame_stride,
+                                                           bool use_time,
                                                            int radius,
                                                            float pca_threshold);
 
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_difference)(int dx,
                                                            int dy,
                                                            float *weight_image,
-                                                           float *variance,
+                                                           float *variance_image,
+                                                           float *scale_image,
                                                            float *difference_image,
                                                            int* rect,
                                                            int stride,
                                                            int channel_offset,
+                                                           int frame_offset,
                                                            float a,
                                                            float k_2);
 
@@ -99,11 +114,13 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_update_output)(int dx,
                                                          float *out_image,
                                                          float *accum_image,
                                                          int* rect,
+                                                         int channel_offset,
                                                          int stride,
                                                          int f);
 
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx,
                                                              int dy,
+                                                             int t,
                                                              float *difference_image,
                                                              float *buffer,
                                                              float *transform,
@@ -114,7 +131,9 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx,
                                                              int *filter_window,
                                                              int stride,
                                                              int f,
-                                                             int pass_stride);
+                                                             int pass_stride,
+                                                             int frame_offset,
+                                                             bool use_time);
 
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_normalize)(float *out_image,
                                                      float *accum_image,
index 4c75871..c295058 100644 (file)
@@ -69,6 +69,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample,
                                                    int x,
                                                    int y,
                                                    float *mean, float *variance,
+                                                   float scale,
                                                    int* prefilter_rect,
                                                    int buffer_pass_stride,
                                                    int buffer_denoising_offset)
@@ -80,12 +81,29 @@ void KERNEL_FUNCTION_FULL_NAME(filter_get_feature)(int sample,
                                  m_offset, v_offset,
                                  x, y,
                                  mean, variance,
+                                 scale,
                                  load_int4(prefilter_rect),
                                  buffer_pass_stride,
                                  buffer_denoising_offset);
 #endif
 }
 
+void KERNEL_FUNCTION_FULL_NAME(filter_write_feature)(int sample,
+                                                     int x,
+                                                     int y,
+                                                     int *buffer_params,
+                                                     float *from,
+                                                     float *buffer,
+                                                     int out_offset,
+                                                     int* prefilter_rect)
+{
+#ifdef KERNEL_STUB
+       STUB_ASSERT(KERNEL_ARCH, filter_write_feature);
+#else
+       kernel_filter_write_feature(sample, x, y, load_int4(buffer_params), from, buffer, out_offset, load_int4(prefilter_rect));
+#endif
+}
+
 void KERNEL_FUNCTION_FULL_NAME(filter_detect_outliers)(int x, int y,
                                                        ccl_global float *image,
                                                        ccl_global float *variance,
@@ -117,6 +135,7 @@ void KERNEL_FUNCTION_FULL_NAME(filter_combine_halves)(int x, int y,
 }
 
 void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer,
+                                                           TileInfo *tile_info,
                                                            int x,
                                                            int y,
                                                            int storage_ofs,
@@ -124,18 +143,23 @@ void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer,
                                                            int *rank,
                                                            int* prefilter_rect,
                                                            int pass_stride,
+                                                           int frame_stride,
+                                                           bool use_time,
                                                            int radius,
                                                            float pca_threshold)
 {
 #ifdef KERNEL_STUB
        STUB_ASSERT(KERNEL_ARCH, filter_construct_transform);
 #else
-  rank += storage_ofs;
-  transform += storage_ofs*TRANSFORM_SIZE;
+       rank += storage_ofs;
+       transform += storage_ofs*TRANSFORM_SIZE;
        kernel_filter_construct_transform(buffer,
+                                         tile_info,
                                          x, y,
                                          load_int4(prefilter_rect),
                                          pass_stride,
+                                         frame_stride,
+                                         use_time,
                                          transform,
                                          rank,
                                          radius,
@@ -146,18 +170,29 @@ void KERNEL_FUNCTION_FULL_NAME(filter_construct_transform)(float* buffer,
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_calc_difference)(int dx,
                                                            int dy,
                                                            float *weight_image,
-                                                           float *variance,
+                                                           float *variance_image,
+                                                           float *scale_image,
                                                            float *difference_image,
                                                            int *rect,
                                                            int stride,
                                                            int channel_offset,
+                                                           int frame_offset,
                                                            float a,
                                                            float k_2)
 {
 #ifdef KERNEL_STUB
        STUB_ASSERT(KERNEL_ARCH, filter_nlm_calc_difference);
 #else
-       kernel_filter_nlm_calc_difference(dx, dy, weight_image, variance, difference_image, load_int4(rect), stride, channel_offset, a, k_2);
+       kernel_filter_nlm_calc_difference(dx, dy,
+                                         weight_image,
+                                         variance_image,
+                                         scale_image,
+                                         difference_image,
+                                         load_int4(rect),
+                                         stride,
+                                         channel_offset,
+                                         frame_offset,
+                                         a, k_2);
 #endif
 }
 
@@ -195,18 +230,28 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_update_output)(int dx,
                                                          float *out_image,
                                                          float *accum_image,
                                                          int *rect,
+                                                         int channel_offset,
                                                          int stride,
                                                          int f)
 {
 #ifdef KERNEL_STUB
        STUB_ASSERT(KERNEL_ARCH, filter_nlm_update_output);
 #else
-       kernel_filter_nlm_update_output(dx, dy, difference_image, image, temp_image, out_image, accum_image, load_int4(rect), stride, f);
+       kernel_filter_nlm_update_output(dx, dy,
+                                       difference_image,
+                                       image,
+                                       temp_image,
+                                       out_image,
+                                       accum_image,
+                                       load_int4(rect),
+                                       channel_offset,
+                                       stride, f);
 #endif
 }
 
 void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx,
                                                              int dy,
+                                                             int t,
                                                              float *difference_image,
                                                              float *buffer,
                                                              float *transform,
@@ -217,12 +262,24 @@ void KERNEL_FUNCTION_FULL_NAME(filter_nlm_construct_gramian)(int dx,
                                                              int *filter_window,
                                                              int stride,
                                                              int f,
-                                                             int pass_stride)
+                                                             int pass_stride,
+                                                             int frame_offset,
+                                                             bool use_time)
 {
 #ifdef KERNEL_STUB
        STUB_ASSERT(KERNEL_ARCH, filter_nlm_construct_gramian);
 #else
-       kernel_filter_nlm_construct_gramian(dx, dy, difference_image, buffer, transform, rank, XtWX, XtWY, load_int4(rect), load_int4(filter_window), stride, f, pass_stride);
+       kernel_filter_nlm_construct_gramian(dx, dy, t,
+                                           difference_image,
+                                           buffer,
+                                           transform, rank,
+                                           XtWX, XtWY,
+                                           load_int4(rect),
+                                           load_int4(filter_window),
+                                           stride, f,
+                                           pass_stride,
+                                           frame_offset,
+                                           use_time);
 #endif
 }
 
index b856cbd..5b552b0 100644 (file)
@@ -29,7 +29,7 @@
 extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
 kernel_cuda_filter_divide_shadow(int sample,
-                                 TileInfo *tile_info,
+                                 CCL_FILTER_TILE_INFO,
                                  float *unfilteredA,
                                  float *unfilteredB,
                                  float *sampleVariance,
@@ -59,11 +59,12 @@ kernel_cuda_filter_divide_shadow(int sample,
 extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
 kernel_cuda_filter_get_feature(int sample,
-                               TileInfo *tile_info,
+                               CCL_FILTER_TILE_INFO,
                                int m_offset,
                                int v_offset,
                                float *mean,
                                float *variance,
+                               float scale,
                                int4 prefilter_rect,
                                int buffer_pass_stride,
                                int buffer_denoising_offset)
@@ -76,12 +77,37 @@ kernel_cuda_filter_get_feature(int sample,
                                          m_offset, v_offset,
                                          x, y,
                                          mean, variance,
+                                         scale,
                                          prefilter_rect,
                                          buffer_pass_stride,
                                          buffer_denoising_offset);
        }
 }
 
+extern "C" __global__ void
+CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
+kernel_cuda_filter_write_feature(int sample,
+                                 int4 buffer_params,
+                                 int4 filter_area,
+                                 float *from,
+                                 float *buffer,
+                                 int out_offset,
+                                 int4 prefilter_rect)
+{
+       int x = blockDim.x*blockIdx.x + threadIdx.x;
+       int y = blockDim.y*blockIdx.y + threadIdx.y;
+       if(x < filter_area.z && y < filter_area.w) {
+               kernel_filter_write_feature(sample,
+                                           x + filter_area.x,
+                                           y + filter_area.y,
+                                           buffer_params,
+                                           from,
+                                           buffer,
+                                           out_offset,
+                                           prefilter_rect);
+       }
+}
+
 extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
 kernel_cuda_filter_detect_outliers(float *image,
@@ -112,10 +138,12 @@ kernel_cuda_filter_combine_halves(float *mean, float *variance, float *a, float
 extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
 kernel_cuda_filter_construct_transform(float const* __restrict__ buffer,
+                                       CCL_FILTER_TILE_INFO,
                                        float *transform, int *rank,
                                        int4 filter_area, int4 rect,
                                        int radius, float pca_threshold,
-                                       int pass_stride)
+                                       int pass_stride, int frame_stride,
+                                       bool use_time)
 {
        int x = blockDim.x*blockIdx.x + threadIdx.x;
        int y = blockDim.y*blockIdx.y + threadIdx.y;
@@ -123,8 +151,11 @@ kernel_cuda_filter_construct_transform(float const* __restrict__ buffer,
                int *l_rank = rank + y*filter_area.z + x;
                float *l_transform = transform + y*filter_area.z + x;
                kernel_filter_construct_transform(buffer,
+                                                 tile_info,
                                                  x + filter_area.x, y + filter_area.y,
-                                                 rect, pass_stride,
+                                                 rect,
+                                                 pass_stride, frame_stride,
+                                                 use_time,
                                                  l_transform, l_rank,
                                                  radius, pca_threshold,
                                                  filter_area.z*filter_area.w,
@@ -136,6 +167,7 @@ extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
 kernel_cuda_filter_nlm_calc_difference(const float *ccl_restrict weight_image,
                                        const float *ccl_restrict variance_image,
+                                       const float *ccl_restrict scale_image,
                                        float *difference_image,
                                        int w,
                                        int h,
@@ -143,6 +175,7 @@ kernel_cuda_filter_nlm_calc_difference(const float *ccl_restrict weight_image,
                                        int pass_stride,
                                        int r,
                                        int channel_offset,
+                                       int frame_offset,
                                        float a,
                                        float k_2)
 {
@@ -152,9 +185,12 @@ kernel_cuda_filter_nlm_calc_difference(const float *ccl_restrict weight_image,
                kernel_filter_nlm_calc_difference(co.x, co.y, co.z, co.w,
                                                  weight_image,
                                                  variance_image,
+                                                 scale_image,
                                                  difference_image + ofs,
                                                  rect, stride,
-                                                 channel_offset, a, k_2);
+                                                 channel_offset,
+                                                 frame_offset,
+                                                 a, k_2);
        }
 }
 
@@ -210,6 +246,7 @@ kernel_cuda_filter_nlm_update_output(const float *ccl_restrict difference_image,
                                      int h,
                                      int stride,
                                      int pass_stride,
+                                     int channel_offset,
                                      int r,
                                      int f)
 {
@@ -221,7 +258,9 @@ kernel_cuda_filter_nlm_update_output(const float *ccl_restrict difference_image,
                                                image,
                                                out_image,
                                                accum_image,
-                                               rect, stride, f);
+                                               rect,
+                                               channel_offset,
+                                               stride, f);
        }
 }
 
@@ -242,7 +281,8 @@ kernel_cuda_filter_nlm_normalize(float *out_image,
 
 extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
-kernel_cuda_filter_nlm_construct_gramian(const float *ccl_restrict difference_image,
+kernel_cuda_filter_nlm_construct_gramian(int t,
+                                         const float *ccl_restrict difference_image,
                                          const float *ccl_restrict buffer,
                                          float const* __restrict__ transform,
                                          int *rank,
@@ -254,13 +294,16 @@ kernel_cuda_filter_nlm_construct_gramian(const float *ccl_restrict difference_im
                                          int stride,
                                          int pass_stride,
                                          int r,
-                                         int f)
+                                         int f,
+                                         int frame_offset,
+                                         bool use_time)
 {
        int4 co, rect;
        int ofs;
        if(get_nlm_coords_window(w, h, r, pass_stride, &rect, &co, &ofs, filter_window)) {
                kernel_filter_nlm_construct_gramian(co.x, co.y,
                                                    co.z, co.w,
+                                                   t,
                                                    difference_image + ofs,
                                                    buffer,
                                                    transform, rank,
@@ -268,6 +311,8 @@ kernel_cuda_filter_nlm_construct_gramian(const float *ccl_restrict difference_im
                                                    rect, filter_window,
                                                    stride, f,
                                                    pass_stride,
+                                                   frame_offset,
+                                                   use_time,
                                                    threadIdx.y*blockDim.x + threadIdx.x);
        }
 }
index a550f97..996bc27 100644 (file)
@@ -56,6 +56,7 @@ __kernel void kernel_ocl_filter_get_feature(int sample,
                                             int v_offset,
                                             ccl_global float *mean,
                                             ccl_global float *variance,
+                                            float scale,
                                             int4 prefilter_rect,
                                             int buffer_pass_stride,
                                             int buffer_denoising_offset)
@@ -68,12 +69,35 @@ __kernel void kernel_ocl_filter_get_feature(int sample,
                                          m_offset, v_offset,
                                          x, y,
                                          mean, variance,
+                                         scale,
                                          prefilter_rect,
                                          buffer_pass_stride,
                                          buffer_denoising_offset);
        }
 }
 
+__kernel void kernel_ocl_filter_write_feature(int sample,
+                                              int4 buffer_params,
+                                              int4 filter_area,
+                                              ccl_global float *from,
+                                              ccl_global float *buffer,
+                                              int out_offset,
+                                              int4 prefilter_rect)
+{
+       int x = get_global_id(0);
+       int y = get_global_id(1);
+       if(x < filter_area.z && y < filter_area.w) {
+               kernel_filter_write_feature(sample,
+                                           x + filter_area.x,
+                                           y + filter_area.y,
+                                           buffer_params,
+                                           from,
+                                           buffer,
+                                           out_offset,
+                                           prefilter_rect);
+       }
+}
+
 __kernel void kernel_ocl_filter_detect_outliers(ccl_global float *image,
                                                 ccl_global float *variance,
                                                 ccl_global float *depth,
@@ -103,11 +127,14 @@ __kernel void kernel_ocl_filter_combine_halves(ccl_global float *mean,
 }
 
 __kernel void kernel_ocl_filter_construct_transform(const ccl_global float *ccl_restrict buffer,
+                                                    CCL_FILTER_TILE_INFO,
                                                     ccl_global float *transform,
                                                     ccl_global int *rank,
                                                     int4 filter_area,
                                                     int4 rect,
                                                     int pass_stride,
+                                                    int frame_stride,
+                                                    char use_time,
                                                     int radius,
                                                     float pca_threshold)
 {
@@ -117,8 +144,11 @@ __kernel void kernel_ocl_filter_construct_transform(const ccl_global float *ccl_
                ccl_global int *l_rank = rank + y*filter_area.z + x;
                ccl_global float *l_transform = transform + y*filter_area.z + x;
                kernel_filter_construct_transform(buffer,
+                                                 CCL_FILTER_TILE_INFO_ARG,
                                                  x + filter_area.x, y + filter_area.y,
-                                                 rect, pass_stride,
+                                                 rect,
+                                                 pass_stride, frame_stride,
+                                                 use_time,
                                                  l_transform, l_rank,
                                                  radius, pca_threshold,
                                                  filter_area.z*filter_area.w,
@@ -128,6 +158,7 @@ __kernel void kernel_ocl_filter_construct_transform(const ccl_global float *ccl_
 
 __kernel void kernel_ocl_filter_nlm_calc_difference(const ccl_global float *ccl_restrict weight_image,
                                                     const ccl_global float *ccl_restrict variance_image,
+                                                    const ccl_global float *ccl_restrict scale_image,
                                                     ccl_global float *difference_image,
                                                     int w,
                                                     int h,
@@ -135,6 +166,7 @@ __kernel void kernel_ocl_filter_nlm_calc_difference(const ccl_global float *ccl_
                                                     int pass_stride,
                                                     int r,
                                                     int channel_offset,
+                                                    int frame_offset,
                                                     float a,
                                                     float k_2)
 {
@@ -144,9 +176,12 @@ __kernel void kernel_ocl_filter_nlm_calc_difference(const ccl_global float *ccl_
                kernel_filter_nlm_calc_difference(co.x, co.y, co.z, co.w,
                                                  weight_image,
                                                  variance_image,
+                                                 scale_image,
                                                  difference_image + ofs,
                                                  rect, stride,
-                                                 channel_offset, a, k_2);
+                                                 channel_offset,
+                                                 frame_offset,
+                                                 a, k_2);
        }
 }
 
@@ -196,6 +231,7 @@ __kernel void kernel_ocl_filter_nlm_update_output(const ccl_global float *ccl_re
                                                   int h,
                                                   int stride,
                                                   int pass_stride,
+                                                  int channel_offset,
                                                   int r,
                                                   int f)
 {
@@ -207,7 +243,9 @@ __kernel void kernel_ocl_filter_nlm_update_output(const ccl_global float *ccl_re
                                                image,
                                                out_image,
                                                accum_image,
-                                               rect, stride, f);
+                                               rect,
+                                               channel_offset,
+                                               stride, f);
        }
 }
 
@@ -224,7 +262,8 @@ __kernel void kernel_ocl_filter_nlm_normalize(ccl_global float *out_image,
        }
 }
 
-__kernel void kernel_ocl_filter_nlm_construct_gramian(const ccl_global float *ccl_restrict difference_image,
+__kernel void kernel_ocl_filter_nlm_construct_gramian(int t,
+                                                      const ccl_global float *ccl_restrict difference_image,
                                                       const ccl_global float *ccl_restrict buffer,
                                                       const ccl_global float *ccl_restrict transform,
                                                       ccl_global int *rank,
@@ -236,13 +275,16 @@ __kernel void kernel_ocl_filter_nlm_construct_gramian(const ccl_global float *cc
                                                       int stride,
                                                       int pass_stride,
                                                       int r,
-                                                      int f)
+                                                      int f,
+                                                      int frame_offset,
+                                                      char use_time)
 {
        int4 co, rect;
        int ofs;
        if(get_nlm_coords_window(w, h, r, pass_stride, &rect, &co, &ofs, filter_window)) {
                kernel_filter_nlm_construct_gramian(co.x, co.y,
                                                    co.z, co.w,
+                                                   t,
                                                    difference_image + ofs,
                                                    buffer,
                                                    transform, rank,
@@ -250,6 +292,8 @@ __kernel void kernel_ocl_filter_nlm_construct_gramian(const ccl_global float *cc
                                                    rect, filter_window,
                                                    stride, f,
                                                    pass_stride,
+                                                   frame_offset,
+                                                   use_time,
                                                    get_local_id(1)*get_local_size(0) + get_local_id(0));
        }
 }
index b4e3c18..0a60fce 100644 (file)
@@ -42,6 +42,7 @@ BufferParams::BufferParams()
 
        denoising_data_pass = false;
        denoising_clean_pass = false;
+       denoising_prefiltered_pass = false;
 
        Pass::add(PASS_COMBINED, passes);
 }
@@ -73,6 +74,7 @@ int BufferParams::get_passes_size()
        if(denoising_data_pass) {
                size += DENOISING_PASS_SIZE_BASE;
                if(denoising_clean_pass) size += DENOISING_PASS_SIZE_CLEAN;
+               if(denoising_prefiltered_pass) size += DENOISING_PASS_SIZE_PREFILTERED;
        }
 
        return align_up(size, 4);
@@ -88,6 +90,20 @@ int BufferParams::get_denoising_offset()
        return offset;
 }
 
+int BufferParams::get_denoising_prefiltered_offset()
+{
+       assert(denoising_prefiltered_pass);
+
+       int offset = get_denoising_offset();
+
+       offset += DENOISING_PASS_SIZE_BASE;
+       if(denoising_clean_pass) {
+               offset += DENOISING_PASS_SIZE_CLEAN;
+       }
+
+       return offset;
+}
+
 /* Render Buffer Task */
 
 RenderTile::RenderTile()
@@ -153,81 +169,62 @@ bool RenderBuffers::get_denoising_pass_rect(int type, float exposure, int sample
                return false;
        }
 
-       float invsample = 1.0f/sample;
-       float scale = invsample;
-       bool variance = (type == DENOISING_PASS_NORMAL_VAR) ||
-                       (type == DENOISING_PASS_ALBEDO_VAR) ||
-                       (type == DENOISING_PASS_DEPTH_VAR) ||
-                       (type == DENOISING_PASS_COLOR_VAR);
+       float scale = 1.0f;
+       float alpha_scale = 1.0f/sample;
+       if(type == DENOISING_PASS_PREFILTERED_COLOR ||
+          type == DENOISING_PASS_CLEAN ||
+          type == DENOISING_PASS_PREFILTERED_INTENSITY) {
+               scale *= exposure;
+       }
+       else if(type == DENOISING_PASS_PREFILTERED_VARIANCE) {
+               scale *= exposure*exposure * (sample - 1);
+       }
 
-       float scale_exposure = scale;
-       if(type == DENOISING_PASS_COLOR || type == DENOISING_PASS_CLEAN) {
-               scale_exposure *= exposure;
+       int offset;
+       if(type == DENOISING_PASS_CLEAN) {
+               /* The clean pass isn't changed by prefiltering, so we use the original one there. */
+               offset = type + params.get_denoising_offset();
        }
-       else if(type == DENOISING_PASS_COLOR_VAR) {
-               scale_exposure *= exposure*exposure;
+       else if (type == DENOISING_PASS_PREFILTERED_COLOR && !params.denoising_prefiltered_pass) {
+               /* If we're not saving the prefiltering result, return the original noisy pass. */
+               offset = params.get_denoising_offset() + DENOISING_PASS_COLOR;
+               scale /= sample;
+       }
+       else {
+               offset = type + params.get_denoising_prefiltered_offset();
        }
 
-       int offset = type + params.get_denoising_offset();
        int pass_stride = params.get_passes_size();
        int size = params.width*params.height;
 
-       if(variance) {
-               /* Approximate variance as E[x^2] - 1/N * (E[x])^2, since online variance
-                * update does not work efficiently with atomics in the kernel. */
-               int mean_offset = offset - components;
-               float *mean = buffer.data() + mean_offset;
-               float *var = buffer.data() + offset;
-               assert(mean_offset >= 0);
-
-               if(components == 1) {
-                       for(int i = 0; i < size; i++, mean += pass_stride, var += pass_stride, pixels++) {
-                               pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale_exposure;
-                       }
+       float *in = buffer.data() + offset;
+
+       if(components == 1) {
+               for(int i = 0; i < size; i++, in += pass_stride, pixels++) {
+                       pixels[0] = in[0]*scale;
                }
-               else if(components == 3) {
-                       for(int i = 0; i < size; i++, mean += pass_stride, var += pass_stride, pixels += 3) {
-                               pixels[0] = max(0.0f, var[0] - mean[0]*mean[0]*invsample)*scale_exposure;
-                               pixels[1] = max(0.0f, var[1] - mean[1]*mean[1]*invsample)*scale_exposure;
-                               pixels[2] = max(0.0f, var[2] - mean[2]*mean[2]*invsample)*scale_exposure;
-                       }
+       }
+       else if(components == 3) {
+               for(int i = 0; i < size; i++, in += pass_stride, pixels += 3) {
+                       pixels[0] = in[0]*scale;
+                       pixels[1] = in[1]*scale;
+                       pixels[2] = in[2]*scale;
                }
-               else {
-                       return false;
+       }
+       else if(components == 4) {
+               /* Since the alpha channel is not involved in denoising, output the Combined alpha channel. */
+               assert(params.passes[0].type == PASS_COMBINED);
+               float *in_combined = buffer.data();
+
+               for(int i = 0; i < size; i++, in += pass_stride, in_combined += pass_stride, pixels += 4) {
+                       pixels[0] = in[0]*scale;
+                       pixels[1] = in[1]*scale;
+                       pixels[2] = in[2]*scale;
+                       pixels[3] = saturate(in_combined[3]*alpha_scale);
                }
        }
        else {
-               float *in = buffer.data() + offset;
-
-               if(components == 1) {
-                       for(int i = 0; i < size; i++, in += pass_stride, pixels++) {
-                               pixels[0] = in[0]*scale_exposure;
-                       }
-               }
-               else if(components == 3) {
-                       for(int i = 0; i < size; i++, in += pass_stride, pixels += 3) {
-                               pixels[0] = in[0]*scale_exposure;
-                               pixels[1] = in[1]*scale_exposure;
-                               pixels[2] = in[2]*scale_exposure;
-                       }
-               }
-               else if(components == 4) {
-                       assert(type == DENOISING_PASS_COLOR);
-
-                       /* Since the alpha channel is not involved in denoising, output the Combined alpha channel. */
-                       assert(params.passes[0].type == PASS_COMBINED);
-                       float *in_combined = buffer.data();
-
-                       for(int i = 0; i < size; i++, in += pass_stride, in_combined += pass_stride, pixels += 4) {
-                               pixels[0] = in[0]*scale_exposure;
-                               pixels[1] = in[1]*scale_exposure;
-                               pixels[2] = in[2]*scale_exposure;
-                               pixels[3] = saturate(in_combined[3]*scale);
-                       }
-               }
-               else {
-                       return false;
-               }
+               return false;
        }
 
        return true;
index 46c3b89..0a01071 100644 (file)
@@ -54,6 +54,10 @@ public:
        bool denoising_data_pass;
        /* If only some light path types should be denoised, an additional pass is needed. */
        bool denoising_clean_pass;
+       /* When we're prefiltering the passes during rendering, we need to keep both the
+        * original and the prefiltered data around because neighboring tiles might still
+        * need the original data. */
+       bool denoising_prefiltered_pass;
 
        /* functions */
        BufferParams();
@@ -63,6 +67,7 @@ public:
        void add_pass(PassType type);
        int get_passes_size();
        int get_denoising_offset();
+       int get_denoising_prefiltered_offset();
 };
 
 /* Render Buffers */
index 355294e..1d1668a 100644 (file)
@@ -286,6 +286,7 @@ NODE_DEFINE(Film)
 
        SOCKET_BOOLEAN(denoising_data_pass,  "Generate Denoising Data Pass",  false);
        SOCKET_BOOLEAN(denoising_clean_pass, "Generate Denoising Clean Pass", false);
+       SOCKET_BOOLEAN(denoising_prefiltered_pass, "Generate Denoising Prefiltered Pass", false);
        SOCKET_INT(denoising_flags, "Denoising Flags", 0);
 
        return type;
@@ -478,6 +479,9 @@ void Film::device_update(Device *device, DeviceScene *dscene, Scene *scene)
                        kfilm->pass_stride += DENOISING_PASS_SIZE_CLEAN;
                        kfilm->use_light_pass = 1;
                }
+               if(denoising_prefiltered_pass) {
+                       kfilm->pass_stride += DENOISING_PASS_SIZE_PREFILTERED;
+               }
        }
 
        kfilm->pass_stride = align_up(kfilm->pass_stride, 4);
index c597db4..8330a4c 100644 (file)
@@ -60,6 +60,7 @@ public:
        vector<Pass> passes;
        bool denoising_data_pass;
        bool denoising_clean_pass;
+       bool denoising_prefiltered_pass;
        int denoising_flags;
        float pass_alpha_threshold;
 
index 3c6e88d..4b8630a 100644 (file)
@@ -689,7 +689,7 @@ DeviceRequestedFeatures Session::get_requested_device_features()
        BakeManager *bake_manager = scene->bake_manager;
        requested_features.use_baking = bake_manager->get_baking();
        requested_features.use_integrator_branched = (scene->integrator->method == Integrator::BRANCHED_PATH);
-       if(params.denoising_passes) {
+       if(params.run_denoising) {
                requested_features.use_denoising = true;
                requested_features.use_shadow_tricks = true;
        }
@@ -927,7 +927,7 @@ void Session::update_status_time(bool show_pause, bool show_done)
                         */
                        substatus += string_printf(", Sample %d/%d", progress.get_current_sample(), num_samples);
                }
-               if(params.use_denoising) {
+               if(params.run_denoising) {
                        substatus += string_printf(", Denoised %d tiles", progress.get_denoised_tiles());
                }
        }
@@ -975,7 +975,7 @@ void Session::render()
        task.requested_tile_size = params.tile_size;
        task.passes_size = tile_manager.params.get_passes_size();
 
-       if(params.use_denoising) {
+       if(params.run_denoising) {
                task.denoising_radius = params.denoising_radius;
                task.denoising_strength = params.denoising_strength;
                task.denoising_feature_strength = params.denoising_feature_strength;
@@ -983,8 +983,13 @@ void Session::render()
 
                assert(!scene->film->need_update);
                task.pass_stride = scene->film->pass_stride;
+               task.target_pass_stride = task.pass_stride;
                task.pass_denoising_data = scene->film->denoising_data_offset;
                task.pass_denoising_clean = scene->film->denoising_clean_offset;
+
+               task.denoising_from_render = true;
+               task.denoising_do_filter = params.full_denoising;
+               task.denoising_write_passes = params.write_denoising_passes;
        }
 
        device->task_add(task);
index c7f5909..cb1d8fe 100644 (file)
@@ -60,8 +60,9 @@ public:
 
        bool display_buffer_linear;
 
-       bool use_denoising;
-       bool denoising_passes;
+       bool run_denoising;
+       bool write_denoising_passes;
+       bool full_denoising;
        int denoising_radius;
        float denoising_strength;
        float denoising_feature_strength;
@@ -94,8 +95,9 @@ public:
 
                use_profiling = false;
 
-               use_denoising = false;
-               denoising_passes = false;
+               run_denoising = false;
+               write_denoising_passes = false;
+               full_denoising = false;
                denoising_radius = 8;
                denoising_strength = 0.0f;
                denoising_feature_strength = 0.0f;