Commit 27a1ff28 authored by qon's avatar qon
Browse files

fix asynchronous GPU processing, synchronization was one step too late before

parent 6532f7f3
...@@ -84,7 +84,6 @@ int bioem_cuda::compareRefMaps(int iProjectionOut, int iConv, const bioem_map& c ...@@ -84,7 +84,6 @@ int bioem_cuda::compareRefMaps(int iProjectionOut, int iConv, const bioem_map& c
if (GPUAsync) if (GPUAsync)
{ {
checkCudaErrors(cudaEventSynchronize(cudaEvent[iConv & 1])); checkCudaErrors(cudaEventSynchronize(cudaEvent[iConv & 1]));
checkCudaErrors(cudaEventRecord(cudaEvent[iConv & 1], cudaStream));
} }
checkCudaErrors(cudaMemcpyAsync(pConvMap_device[iConv & 1], &conv_map, sizeof(bioem_map), cudaMemcpyHostToDevice, cudaStream)); checkCudaErrors(cudaMemcpyAsync(pConvMap_device[iConv & 1], &conv_map, sizeof(bioem_map), cudaMemcpyHostToDevice, cudaStream));
...@@ -133,7 +132,14 @@ int bioem_cuda::compareRefMaps(int iProjectionOut, int iConv, const bioem_map& c ...@@ -133,7 +132,14 @@ int bioem_cuda::compareRefMaps(int iProjectionOut, int iConv, const bioem_map& c
cout << "Invalid GPU Algorithm selected\n"; cout << "Invalid GPU Algorithm selected\n";
exit(1); exit(1);
} }
if (GPUAsync == 0) checkCudaErrors(cudaStreamSynchronize(cudaStream)); if (GPUAsync)
{
checkCudaErrors(cudaEventRecord(cudaEvent[iConv & 1], cudaStream));
}
else
{
checkCudaErrors(cudaStreamSynchronize(cudaStream));
}
return(0); return(0);
} }
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment