summaryrefslogtreecommitdiffstats
path: root/cuda/3d/cone_bp.cu
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2021-11-16 13:44:13 +0100
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2021-11-16 14:06:30 +0100
commit39582115bc93b5435d25e56891815ae7cb1898fd (patch)
tree6c728e1125961fc04ba6f77bf9af637925825f5b /cuda/3d/cone_bp.cu
parentb492e3d049e300132d2f22eee7922ff308342a84 (diff)
downloadastra-39582115bc93b5435d25e56891815ae7cb1898fd.tar.gz
astra-39582115bc93b5435d25e56891815ae7cb1898fd.tar.bz2
astra-39582115bc93b5435d25e56891815ae7cb1898fd.tar.xz
astra-39582115bc93b5435d25e56891815ae7cb1898fd.zip
Remove cudaTextForceKernelsCompletion
Diffstat (limited to 'cuda/3d/cone_bp.cu')
-rw-r--r--cuda/3d/cone_bp.cu4
1 files changed, 3 insertions, 1 deletions
diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index 7c3fc8d..e265304 100644
--- a/cuda/3d/cone_bp.cu
+++ b/cuda/3d/cone_bp.cu
@@ -357,7 +357,9 @@ bool ConeBP_Array(cudaPitchedPtr D_volumeData,
dev_cone_BP_SS<<<dimGrid, dimBlock>>>(D_volumeData.ptr, D_volumeData.pitch/sizeof(float), i, th, dims, params.iRaysPerVoxelDim, fOutputScale);
}
- cudaTextForceKernelsCompletion();
+ // TODO: Consider not synchronizing here, if possible.
+ if (!checkCuda(cudaThreadSynchronize(), "cone_bp"))
+ return false;
angles = angles + angleCount;
// printf("%f\n", toc(t));