Fix artifacts in t-SNE results (#3084)

Fixes #3057 Co-authored-by: Corey J. Nolet <[email protected]>
rapidsai · Nov 2, 2020 · 898f480 · 898f480
1 parent 5b7757a
commit 898f480
Show file tree

Hide file tree

Showing 3 changed files with 4 additions and 3 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -33,6 +33,7 @@
 - PR #3061: Handle C++ exception thrown from FIL predict
 - PR #3073: Update mathjax CDN URL for documentation
 - PR #3062: Bumping xgboost version to match cuml version
+- PR #3084: Fix artifacts in t-SNE results
 - PR #3086: Reverting FIL Notebook Testing
 
 # cuML 0.16.0 (Date TBD)

diff --git a/cpp/src/tsne/barnes_hut.cuh b/cpp/src/tsne/barnes_hut.cuh
@@ -149,8 +149,8 @@ void Barnes_Hut(float *VAL, const int *COL, const int *ROW, const int NNZ,
     cudaFuncSetCacheConfig(TSNE::TreeBuildingKernel, cudaFuncCachePreferL1));
   CUDA_CHECK(cudaFuncSetCacheConfig(TSNE::ClearKernel1, cudaFuncCachePreferL1));
   CUDA_CHECK(cudaFuncSetCacheConfig(TSNE::ClearKernel2, cudaFuncCachePreferL1));
-  CUDA_CHECK(
-    cudaFuncSetCacheConfig(TSNE::SummarizationKernel, cudaFuncCachePreferL1));
+  CUDA_CHECK(cudaFuncSetCacheConfig(TSNE::SummarizationKernel,
+                                    cudaFuncCachePreferShared));
   CUDA_CHECK(cudaFuncSetCacheConfig(TSNE::SortKernel, cudaFuncCachePreferL1));
   CUDA_CHECK(
     cudaFuncSetCacheConfig(TSNE::RepulsionKernel, cudaFuncCachePreferL1));

diff --git a/cpp/src/tsne/bh_kernels.cuh b/cpp/src/tsne/bh_kernels.cuh
@@ -449,7 +449,7 @@ __global__ __launch_bounds__(THREADS3, FACTOR3) void SummarizationKernel(
     }
 
   SKIP_LOOP:
-    __syncthreads();
+    __threadfence();
     if (flag != 0) {
       massd[k] = cm;
       k += inc;