diff options
author | Patrick Mours <pmours@nvidia.com> | 2020-05-27 16:31:03 +0300 |
---|---|---|
committer | Patrick Mours <pmours@nvidia.com> | 2020-05-27 16:31:03 +0300 |
commit | 49c295813ba9fd9de3209e50a11e55647c1de654 (patch) | |
tree | 2b006406199642d0ed666dcac98b2ea138c76f18 /intern | |
parent | efa4ae17f043f41b2d7d26dc8dfc7a7f90b4acfd (diff) | |
parent | 28d9368538f93d329573910d080e2f753f91b489 (diff) |
Merge branch 'blender-v2.83-release'
Diffstat (limited to 'intern')
-rw-r--r-- | intern/cycles/device/device_optix.cpp | 3 | ||||
-rw-r--r-- | intern/cycles/kernel/kernels/cuda/filter.cu | 14 |
2 files changed, 9 insertions, 8 deletions
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; } } |