[Bf-blender-cvs] [aa5426d96ba] cycles-x: Fix OptiX denoiser for transparent film in Cycles X

[Sergey Sharybin]

Commit: aa5426d96bac0335e15c48b961dbf3b405cfec3a
Author: Sergey Sharybin
Date:   Tue Jun 15 14:44:19 2021 +0200
Branches: cycles-x
https://developer.blender.org/rBaa5426d96bac0335e15c48b961dbf3b405cfec3a

Fix OptiX denoiser for transparent film in Cycles X

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

M	intern/cycles/device/optix/device_impl.cpp
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 b9f624f0ba0..46369e18f49 100644
--- a/intern/cycles/device/optix/device_impl.cpp
+++ b/intern/cycles/device/optix/device_impl.cpp
@@ -616,8 +616,8 @@ bool OptiXDevice::denoise_filter_convert_from_rgb(OptiXDeviceQueue *queue,
 {
   const int work_size = task.buffer_params.width * task.buffer_params.height;
 
-  const int pass_combined_denoised = task.buffer_params.get_pass_offset(PASS_COMBINED,
-                                                                        PassMode::DENOISED);
+  const int pass_noisy = task.buffer_params.get_pass_offset(PASS_COMBINED, PassMode::NOISY);
+  const int pass_denoised = task.buffer_params.get_pass_offset(PASS_COMBINED, PassMode::DENOISED);
   const int pass_sample_count = task.buffer_params.get_pass_offset(PASS_SAMPLE_COUNT);
 
   void *args[] = {const_cast<device_ptr *>(&d_input_rgb),
@@ -630,7 +630,8 @@ bool OptiXDevice::denoise_filter_convert_from_rgb(OptiXDeviceQueue *queue,
                   const_cast<int *>(&task.buffer_params.stride),
                   const_cast<int *>(&task.buffer_params.pass_stride),
                   const_cast<int *>(&task.num_samples),
-                  const_cast<int *>(&pass_combined_denoised),
+                  const_cast<int *>(&pass_noisy),
+                  const_cast<int *>(&pass_denoised),
                   const_cast<int *>(&pass_sample_count)};
 
   return queue->enqueue(DEVICE_KERNEL_FILTER_CONVERT_FROM_RGB, work_size, args);
diff --git a/intern/cycles/kernel/device/cuda/kernel.cu b/intern/cycles/kernel/device/cuda/kernel.cu
index 13f89c7d9c1..3cedaa3623a 100644
--- a/intern/cycles/kernel/device/cuda/kernel.cu
+++ b/intern/cycles/kernel/device/cuda/kernel.cu
@@ -681,6 +681,7 @@ extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_KERNEL_BLOCK_NUM_THREADS,
                                         int stride,
                                         int pass_stride,
                                         int num_samples,
+                                        int pass_noisy,
                                         int pass_denoised,
                                         int pass_sample_count)
 {
@@ -705,11 +706,13 @@ extern "C" __global__ void CUDA_LAUNCH_BOUNDS(CUDA_KERNEL_BLOCK_NUM_THREADS,
     pixel_scale = __float_as_uint(buffer[pass_sample_count]);
   }
 
+  const float *noisy_pixel = buffer + pass_noisy;
   float *denoised_pixel = buffer + pass_denoised;
 
   denoised_pixel[0] = in[0] * pixel_scale;
   denoised_pixel[1] = in[1] * pixel_scale;
   denoised_pixel[2] = in[2] * pixel_scale;
+  denoised_pixel[3] = noisy_pixel[3];
 }
 
 #endif