                                             dev_accumulate_weights,
                                             make_int3(input_vol_dim[0], input_vol_dim[1], input_vol_dim[2]),
                                             make_int3(output_vol_dim[0], output_vol_dim[1], output_vol_dim[2]));
+  cudaDeviceSynchronize();
+  CUDA_CHECK_ERROR;
 
   normalize_3Dgrid <<< dimGrid, dimBlock >>> ( dev_output_vol,
-                                            dev_accumulate_weights,
-                                            make_int3(output_vol_dim[0], output_vol_dim[1], output_vol_dim[2]));
-
+                                               dev_accumulate_weights,
+                                               make_int3(output_vol_dim[0], output_vol_dim[1], output_vol_dim[2]));
+  cudaDeviceSynchronize();
+  CUDA_CHECK_ERROR;
+
+  fillHoles_3Dgrid <<< dimGrid, dimBlock >>> ( dev_output_vol,
+                                               dev_accumulate_weights,
+                                               make_int3(output_vol_dim[0], output_vol_dim[1], output_vol_dim[2]));
+  cudaDeviceSynchronize();
   CUDA_CHECK_ERROR;
 
   // Unbind the image and projection matrix textures

//_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_-_

__global__
void normalize_3Dgrid(float * dev_vol_out, float * dev_accumulate_weights, int3 out_dim)
{
  unsigned int i = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;

  cudaMalloc3DArray((cudaArray**)&array_zdvf, &channelDesc, dvfExtent);
  CUDA_CHECK_ERROR;

  // Copy image data to arrays. For a nice explanation on make_cudaPitchedPtr, checkout
  // http://stackoverflow.com/questions/16119943/how-and-when-should-i-use-pitched-pointer-with-the-cuda-api
  cudaMemcpy3DParms xCopyParams = {0};
  xCopyParams.srcPtr   = make_cudaPitchedPtr(dev_input_xdvf, input_dvf_dim[0] * sizeof(float), input_dvf_dim[0], input_dvf_dim[1]);
