fix synchronization

Author: mgrabban

tsne/CUDA/src/kernels/perplexity_search.cu

@@ -153,12 +153,12 @@ void tsnecuda::SearchPerplexity(
 
         // compute entropy of current row
         row_sum = tsnecuda::utils::ReduceSum(handle, pij, num_neighbors, num_points, 0);
-        GpuErrorCheck(cudaDeviceSynchronize());
+        // GpuErrorCheck(cudaDeviceSynchronize());
 
         // compute negative entropy
         thrust::transform(pij.begin(), pij.end(), entropy.begin(), tsnecuda::utils::FunctionalEntropy());
         neg_entropy = tsnecuda::utils::ReduceAlpha(handle, entropy, num_neighbors, num_points, -1.0f, 0);
-        GpuErrorCheck(cudaDeviceSynchronize());
+        // GpuErrorCheck(cudaDeviceSynchronize());
 
         // binary search for beta
         PerplexitySearchKernel<<<NBLOCKS2, BLOCKSIZE2>>>(

tsne/CUDA/src/utils/reduce_utils.cu

@@ -57,6 +57,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(cudaDeviceSynchronize());
         return means;
     } else if (axis == 1) {
         thrust::device_vector<float> ones(M, 1.f);
@@ -67,6 +68,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(cudaDeviceSynchronize());
         return means;
     } else {
         throw std::runtime_error("Axis must be 0 or 1.");

tsne/HIP/src/kernels/perplexity_search.cpp

@@ -154,12 +154,12 @@ void tsnecuda::SearchPerplexity(
 
         // compute entropy of current row
         row_sum = tsnecuda::utils::ReduceSum(handle, pij, num_neighbors, num_points, 0);
-        GpuErrorCheck(hipDeviceSynchronize());
+        // GpuErrorCheck(hipDeviceSynchronize());
 
         // compute negative entropy
         thrust::transform(pij.begin(), pij.end(), entropy.begin(), tsnecuda::utils::FunctionalEntropy());
         neg_entropy = tsnecuda::utils::ReduceAlpha(handle, entropy, num_neighbors, num_points, -1.0f, 0);
-        GpuErrorCheck(hipDeviceSynchronize());
+        // GpuErrorCheck(hipDeviceSynchronize());
 
         // binary search for beta
         hipLaunchKernelGGL(PerplexitySearchKernel, NBLOCKS2, BLOCKSIZE2, 0, 0, 

tsne/HIP/src/utils/reduce_utils.cpp

@@ -57,6 +57,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(hipDeviceSynchronize());
         return means;
     } else if (axis == 1) {
         thrust::device_vector<float> ones(M, 1.f);
@@ -67,6 +68,7 @@ thrust::device_vector<float> tsnecuda::utils::ReduceAlpha(
                 thrust::raw_pointer_cast(d_matrix.data()), N,
                 thrust::raw_pointer_cast(ones.data()), 1, &kBeta,
                 thrust::raw_pointer_cast(means.data()), 1));
+        GpuErrorCheck(hipDeviceSynchronize());
         return means;
     } else {
         throw std::runtime_error("Axis must be 0 or 1.");