Commit: 8478bce73b537aec5ba3a3428e628374822598c9
Author: Lukas Stockner
Date: Wed Nov 16 16:10:29 2016 +0100
Branches: soc-2016-cycles_denoising
https://developer.blender.org/rB8478bce73b537aec5ba3a3428e628374822598c9
Cycles: Fix Collaborative CUDA WLR denoising
===================================================================
M intern/cycles/device/device_cuda.cpp
===================================================================
diff --git a/intern/cycles/device/device_cuda.cpp
b/intern/cycles/device/device_cuda.cpp
index ec7a4e7..e721118 100644
--- a/intern/cycles/device/device_cuda.cpp
+++ b/intern/cycles/device/device_cuda.cpp
@@ -1132,16 +1132,6 @@ public:
0, 0, final_args, 0));
cuda_assert(cuCtxSynchronize());
-
- void *divide_args[] = {&d_buffers,
- &sample,
- &rtile.offset,
- &rtile.stride,
- &filter_area};
- cuda_assert(cuLaunchKernel(cuFilterDivideCombined,
- xblocks , yblocks, 1, /*
blocks */
- xthreads, ythreads, 1, /*
threads */
- 0, 0, divide_args, 0));
}
else {
cuda_assert(cuLaunchKernel(cuFilterEstimateBandwidths,
@@ -1188,6 +1178,18 @@ public:
cuda_assert(cuCtxSynchronize());
}
+ if(kernel_globals.integrator.use_collaborative_filtering) {
+ void *divide_args[] = {&d_buffers,
+ &sample,
+ &rtile.offset,
+ &rtile.stride,
+ &filter_area};
+ cuda_assert(cuLaunchKernel(cuFilterDivideCombined,
+ xblocks , yblocks, 1, /*
blocks */
+ xthreads, ythreads, 1, /*
threads */
+ 0, 0, divide_args, 0));
+ }
+
#ifdef WITH_CYCLES_DEBUG_FILTER
CUDAFilterStorage *host_storage = new
CUDAFilterStorage[filter_area.z*filter_area.w];
cuda_assert(cuMemcpyDtoH(host_storage, d_storage,
sizeof(CUDAFilterStorage)*filter_area.z*filter_area.w));
_______________________________________________
Bf-blender-cvs mailing list
[email protected]
https://lists.blender.org/mailman/listinfo/bf-blender-cvs
