Cr8-xyz/blender
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Merge branch 'blender2.7'

Browse Source
This commit is contained in:
Brecht Van Lommel
2019-02-06 15:22:53 +01:00
parent e8292466bc fccf506ed7
commit e21ae0bb26
31 changed files with 1047 additions and 417 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
intern/cycles/blender/addon/engine.py
View 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",

17
intern/cycles/blender/blender_session.cpp
View 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");

35
intern/cycles/blender/blender_sync.cpp
View 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());

90
intern/cycles/device/device_cpu.cpp
View 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);

174
intern/cycles/device/device_cuda.cpp
View 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"));
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_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);
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};
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_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);

119
intern/cycles/device/device_denoising.cpp
View 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();
}
construct_transform();
reconstruct();
if(do_filter) {
construct_transform();
reconstruct();
}
if(write_passes) {
write_buffer();
}
functions.unmap_neighbor_tiles(rtiles);
}

27
intern/cycles/device/device_denoising.h
View 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

8
intern/cycles/device/device_task.h
View 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;

16
intern/cycles/device/opencl/opencl.h
View 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,

100
intern/cycles/device/opencl/opencl_base.cpp
View 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);

7
intern/cycles/kernel/filter/filter_defines.h
View File

@@ -17,16 +17,21 @@
#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];

77
intern/cycles/kernel/filter/filter_features.h
View File

@@ -18,19 +18,23 @@
#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

57
intern/cycles/kernel/filter/filter_features_sse.h
View 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(mean) {
for(int i = 0; i < DENOISE_FEATURES; i++)
features[i] = features[i] - mean[i];
if(use_time) {
features[10] = t;
}
for(int i = 0; i < DENOISE_FEATURES; i++)
if(mean) {
for(int i = 0; i < num_features; i++) {
features[i] = features[i] - mean[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)));

32
intern/cycles/kernel/filter/filter_nlm_cpu.h
View 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);

45
intern/cycles/kernel/filter/filter_nlm_gpu.h
View 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,

82
intern/cycles/kernel/filter/filter_prefilter.h
View 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;

24
intern/cycles/kernel/filter/filter_reconstruction.h
View 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;

59
intern/cycles/kernel/filter/filter_transform.h
View File

@@ -17,8 +17,10 @@
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

54
intern/cycles/kernel/filter/filter_transform_gpu.h
View File

@@ -17,8 +17,10 @@
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];
}
}
}

60
intern/cycles/kernel/filter/filter_transform_sse.h
View File

@@ -17,8 +17,10 @@
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);
}
}

9
intern/cycles/kernel/kernel_types.h
View 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 {

23
intern/cycles/kernel/kernels/cpu/filter_cpu.h
View 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,

71
intern/cycles/kernel/kernels/cpu/filter_cpu_impl.h
View 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
}

61
intern/cycles/kernel/kernels/cuda/filter.cu
View 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);
}
}

54
intern/cycles/kernel/kernels/opencl/filter.cl
View 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));
}
}

123
intern/cycles/render/buffers.cpp
View 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_exposure = scale;
if(type == DENOISING_PASS_COLOR || type == DENOISING_PASS_CLEAN) {
scale_exposure *= exposure;
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_COLOR_VAR) {
scale_exposure *= exposure*exposure;
else if(type == DENOISING_PASS_PREFILTERED_VARIANCE) {
scale *= exposure*exposure * (sample - 1);
}
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_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);
float *in = buffer.data() + offset;
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;
}
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;

5
intern/cycles/render/buffers.h
View 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 */

4
intern/cycles/render/film.cpp
View 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);

1
intern/cycles/render/film.h
View 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;

11
intern/cycles/render/session.cpp
View 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);

10
intern/cycles/render/session.h
View 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;