added error handling for bspline z-dir

Author: krzysg

src/algorithm/bsplineZdir.cuh

@@ -129,24 +129,15 @@ __global__ void bsplineZdir(T *image, PixelDataDim dim, BsplineParamsCuda p, boo
  * Function for launching a kernel
  */
 template<typename T>
-void runBsplineZdir(T *cudaImage, PixelDataDim dim, BsplineParamsCuda &p, cudaStream_t aStream) {
+void runBsplineZdir(T *cudaImage, PixelDataDim dim, BsplineParamsCuda &p, bool *error, cudaStream_t aStream) {
     constexpr int numOfWorkersYdir = 128;
     dim3 threadsPerBlockZ(1, numOfWorkersYdir, 1);
     dim3 numBlocksZ(1,
                     (dim.y + threadsPerBlockZ.y - 1) / threadsPerBlockZ.y,
                     (dim.x + threadsPerBlockZ.x - 1) / threadsPerBlockZ.x);
     // In case of error this will be set to true by one of the kernels (CUDA does not guarantee which kernel will set global variable if more then one kernel
     // access it but this is enough for us to know that somewhere in one on more kernels overflow was detected.
-    bool isErrorDetected = false;
-    {
-        ScopedCudaMemHandler<bool*, H2D | D2H> error(&isErrorDetected, 1, aStream);
-        bsplineZdir<T> <<<numBlocksZ, threadsPerBlockZ, 0, aStream>>> (cudaImage, dim, p, error.get());
-    }
-
-    if (isErrorDetected) {
-        throw std::invalid_argument("integer under-/overflow encountered in CUDA bsplineZdir - "
-                                    "try squashing the input image to a narrower range or use APRConverter<float>");
-    }
+    bsplineZdir<T> <<<numBlocksZ, threadsPerBlockZ, 0, aStream>>> (cudaImage, dim, p, error);
 }
 
 #endif