[Bf-blender-cvs] [15656d1f36d] cycles-x: Cycles X: Use pass accessor in the OptiX denoiser

[Sergey Sharybin]

Commit: 15656d1f36d3f1ed4b9cc68076559a4bce2f16ee
Author: Sergey Sharybin
Date:   Tue Jun 15 14:35:10 2021 +0200
Branches: cycles-x
https://developer.blender.org/rB15656d1f36d3f1ed4b9cc68076559a4bce2f16ee

Cycles X: Use pass accessor in the OptiX denoiser

No functional changes, just converging pass accessor to a commonly
used utility.

There is still custom kernel used, which ensures that pixel values
are within expected range.

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

M	intern/cycles/device/optix/device_impl.cpp
M	intern/cycles/device/optix/device_impl.h
M	intern/cycles/kernel/device/cuda/kernel.cu

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

diff --git a/intern/cycles/device/optix/device_impl.cpp b/intern/cycles/device/optix/device_impl.cpp
index a5fccb92c08..ef566cff8d3 100644
--- a/intern/cycles/device/optix/device_impl.cpp
+++ b/intern/cycles/device/optix/device_impl.cpp
@@ -21,6 +21,7 @@
 
 #  include "bvh/bvh.h"
 #  include "bvh/bvh_optix.h"
+#  include "integrator/pass_accessor_gpu.h"
 #  include "render/buffers.h"
 #  include "render/hair.h"
 #  include "render/mesh.h"
@@ -553,10 +554,13 @@ void OptiXDevice::denoise_buffer(const DeviceDenoiseTask &task)
 
   OptiXDeviceQueue queue(this);
 
+  /* Read pixels from the noisy input pass, store them in the temporary buffer for further
+   * clamping. */
+  denoise_read_input_pixels(&queue, task, input_rgb.device_pointer);
+
   /* Make sure input data is in [0 .. 10000] range by scaling the input buffer by the number of
-   *
-   * samples in the buffer. This will do (scaled) copy of the noisy image and needed passes into
-   * an input buffer for the OptiX denoiser. */
+   * samples in the buffer. Additionally, fill in the auxillary passes needed by the denoiser which
+   * were not provided by the pass accessor. */
   if (!denoise_filter_convert_to_rgb(&queue, task, input_rgb.device_pointer)) {
     LOG(ERROR) << "Error connverting denoising passes to RGB buffer.";
     return;
@@ -579,6 +583,33 @@ void OptiXDevice::denoise_buffer(const DeviceDenoiseTask &task)
   queue.synchronize();
 }
 
+void OptiXDevice::denoise_read_input_pixels(OptiXDeviceQueue *queue,
+                                            const DeviceDenoiseTask &task,
+                                            const device_ptr d_input_rgb) const
+{
+  PassAccessor::PassAccessInfo pass_access_info;
+  pass_access_info.type = PASS_COMBINED;
+  pass_access_info.mode = PassMode::NOISY;
+  pass_access_info.offset = task.buffer_params.get_pass_offset(pass_access_info.type,
+                                                               pass_access_info.mode);
+
+  /* Denoiser operates on passes which are used to calculate the approximation, and is never used
+   * on the approximation. The latter is not even possible because OptiX does not support
+   * denoising of semi-transparent pixels. */
+  pass_access_info.use_approximate_shadow_catcher = false;
+  pass_access_info.show_active_pixels = false;
+
+  /* TODO(sergey): Consider adding support of actual exposure, to avoid clamping in extreme cases.
+   */
+  const PassAccessorGPU pass_accessor(queue, pass_access_info, 1.0f, task.num_samples);
+
+  PassAccessor::Destination destination(pass_access_info.type);
+  destination.d_pixels = d_input_rgb;
+  destination.num_components = 3;
+
+  pass_accessor.get_render_tile_pixels(task.render_buffers, task.buffer_params, destination);
+}
+
 bool OptiXDevice::denoise_filter_convert_to_rgb(OptiXDeviceQueue *queue,
                                                 const DeviceDenoiseTask &task,
                                                 const device_ptr d_input_rgb)
diff --git a/intern/cycles/device/optix/device_impl.h b/intern/cycles/device/optix/device_impl.h
index 4f3a3a3effe..2bf9516f943 100644
--- a/intern/cycles/device/optix/device_impl.h
+++ b/intern/cycles/device/optix/device_impl.h
@@ -141,6 +141,11 @@ class OptiXDevice : public CUDADevice {
 
   virtual void denoise_buffer(const DeviceDenoiseTask &task) override;
 
+  /* Read pixels from the input noisy image and store scaled result in the given memory. */
+  void denoise_read_input_pixels(OptiXDeviceQueue *queue,
+                                 const DeviceDenoiseTask &task,
+                                 const device_ptr d_input_rgb) const;
+
   /* Run corresponding conversion kernels, preparing data for the denoiser or copying data from the
    * denoiser result to the render buffer. */
   bool denoise_filter_convert_to_rgb(OptiXDeviceQueue *queue,
diff --git a/intern/cycles/kernel/device/cuda/kernel.cu b/intern/cycles/kernel/device/cuda/kernel.cu
index 5327c64893f..e22f53d324d 100644
--- a/intern/cycles/kernel/device/cuda/kernel.cu
+++ b/intern/cycles/kernel/device/cuda/kernel.cu
@@ -683,11 +683,10 @@ extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_KERNEL_BLOCK_NUM_THREADS,
   }
 
   if (num_inputs > 0) {
-    const float *in = buffer + pass_offset.x;
     float *out = rgb + (x + y * sw) * 3;
-    out[0] = clamp(in[0] * pixel_scale, 0.0f, 10000.0f);
-    out[1] = clamp(in[1] * pixel_scale, 0.0f, 10000.0f);
-    out[2] = clamp(in[2] * pixel_scale, 0.0f, 10000.0f);
+    out[0] = clamp(out[0], 0.0f, 10000.0f);
+    out[1] = clamp(out[1], 0.0f, 10000.0f);
+    out[2] = clamp(out[2], 0.0f, 10000.0f);
   }
 
   if (num_inputs > 1) {