[Bf-blender-cvs] [28d9368538f] blender-v2.83-release: Fix T76947: Optix realtime denoiser progressively reduces brightness of very bright objects

[Patrick Mours]

Commit: 28d9368538f93d329573910d080e2f753f91b489
Author: Patrick Mours
Date:   Wed May 27 15:17:47 2020 +0200
Branches: blender-v2.83-release
https://developer.blender.org/rB28d9368538f93d329573910d080e2f753f91b489

Fix T76947: Optix realtime denoiser progressively reduces brightness of very bright objects

The input data to the OptiX denoiser was clamped to 0..10000 as required, but it could easily
exceed that range with a high number of samples (since the data contains the overall sum). To
fix that, divide by the number of samples first and multiply it back in after the denoiser ran.

===================================================================

M	intern/cycles/device/device_optix.cpp
M	intern/cycles/kernel/kernels/cuda/filter.cu

===================================================================

diff --git a/intern/cycles/device/device_optix.cpp b/intern/cycles/device/device_optix.cpp
index e839d852127..441fa35f8af 100644
--- a/intern/cycles/device/device_optix.cpp
+++ b/intern/cycles/device/device_optix.cpp
@@ -924,7 +924,8 @@ class OptiXDevice : public CUDADevice {
                              &rtiles[9].h,
                              &rtiles[9].offset,
                              &rtiles[9].stride,
-                             &task.pass_stride};
+                             &task.pass_stride,
+                             &rtile.sample};
       launch_filter_kernel(
           "kernel_cuda_filter_convert_from_rgb", rtiles[9].w, rtiles[9].h, output_args);
 #  endif
diff --git a/intern/cycles/kernel/kernels/cuda/filter.cu b/intern/cycles/kernel/kernels/cuda/filter.cu
index 22fd5ea5634..6c9642d1f03 100644
--- a/intern/cycles/kernel/kernels/cuda/filter.cu
+++ b/intern/cycles/kernel/kernels/cuda/filter.cu
@@ -57,9 +57,9 @@ kernel_cuda_filter_convert_to_rgb(float *rgb, float *buf, int sw, int sh, int st
 		if (num_inputs > 0) {
 			float *in = buf + x * pass_stride + (y * stride + pass_offset.x) / sizeof(float);
 			float *out = rgb + (x + y * sw) * 3;
-			out[0] = clamp(in[0], 0.0f, 10000.0f);
-			out[1] = clamp(in[1], 0.0f, 10000.0f);
-			out[2] = clamp(in[2], 0.0f, 10000.0f);
+			out[0] = clamp(in[0] / num_samples, 0.0f, 10000.0f);
+			out[1] = clamp(in[1] / num_samples, 0.0f, 10000.0f);
+			out[2] = clamp(in[2] / num_samples, 0.0f, 10000.0f);
 		}
 		if (num_inputs > 1) {
 			float *in = buf + x * pass_stride + (y * stride + pass_offset.y) / sizeof(float);
@@ -80,16 +80,16 @@ kernel_cuda_filter_convert_to_rgb(float *rgb, float *buf, int sw, int sh, int st
 
 extern "C" __global__ void
 CUDA_LAUNCH_BOUNDS(CUDA_THREADS_BLOCK_WIDTH, CUDA_KERNEL_MAX_REGISTERS)
-kernel_cuda_filter_convert_from_rgb(float *rgb, float *buf, int ix, int iy, int iw, int ih, int sx, int sy, int sw, int sh, int offset, int stride, int pass_stride)
+kernel_cuda_filter_convert_from_rgb(float *rgb, float *buf, int ix, int iy, int iw, int ih, int sx, int sy, int sw, int sh, int offset, int stride, int pass_stride, int num_samples)
 {
 	int x = blockDim.x*blockIdx.x + threadIdx.x;
 	int y = blockDim.y*blockIdx.y + threadIdx.y;
 	if(x < sw && y < sh) {
 		float *in = rgb + ((ix + x) + (iy + y) * iw) * 3;
 		float *out = buf + (offset + (sx + x) + (sy + y) * stride) * pass_stride;
-		out[0] = in[0];
-		out[1] = in[1];
-		out[2] = in[2];
+		out[0] = in[0] * num_samples;
+		out[1] = in[1] * num_samples;
+		out[2] = in[2] * num_samples;
 	}
 }