From 39582115bc93b5435d25e56891815ae7cb1898fd Mon Sep 17 00:00:00 2001
From: Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>
Date: Tue, 16 Nov 2021 13:44:13 +0100
Subject: Remove cudaTextForceKernelsCompletion

---
 cuda/2d/par_bp.cu | 10 +++-------
 1 file changed, 3 insertions(+), 7 deletions(-)

(limited to 'cuda/2d/par_bp.cu')

diff --git a/cuda/2d/par_bp.cu b/cuda/2d/par_bp.cu
index b50b5a5..d7c3ab0 100644
--- a/cuda/2d/par_bp.cu
+++ b/cuda/2d/par_bp.cu
@@ -231,13 +231,12 @@ bool BP_internal(float* D_volumeData, unsigned int volumePitch,
 		else
 			devBP<<<dimGrid, dimBlock, 0, stream>>>(D_volumeData, volumePitch, i, dims, fOutputScale);
 	}
-	cudaThreadSynchronize();
 
-	cudaTextForceKernelsCompletion();
+	bool ok = checkCuda(cudaStreamSynchronize(stream), "par_bp");
 
 	cudaStreamDestroy(stream);
 
-	return true;
+	return ok;
 }
 
 bool BP(float* D_volumeData, unsigned int volumePitch,
@@ -284,11 +283,8 @@ bool BP_SART(float* D_volumeData, unsigned int volumePitch,
 	             (dims.iVolHeight+g_blockSliceSize-1)/g_blockSliceSize);
 
 	devBP_SART<<<dimGrid, dimBlock>>>(D_volumeData, volumePitch, angle_offset, angle_scaled_sin, angle_scaled_cos, dims, fOutputScale);
-	cudaThreadSynchronize();
-
-	cudaTextForceKernelsCompletion();
 
-	return true;
+	return checkCuda(cudaThreadSynchronize(), "BP_SART");
 }
 
 
-- 
cgit v1.2.3