Skip to content

Commit

Permalink
fix race in fusedL2knn smem read/write by adding a syncwarp (#679)
Browse files Browse the repository at this point in the history
-- fix race in fusedL2knn smem read/write by adding a syncwarp reported by racecheck tool of compute-sanitizer.
-- this addresses issue - #676

Authors:
  - Mahesh Doijade (https://github.com/mdoijade)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #679
  • Loading branch information
mdoijade authored May 25, 2022
1 parent 372a9a5 commit 29b3a18
Showing 1 changed file with 1 addition and 0 deletions.
1 change: 1 addition & 0 deletions cpp/include/raft/spatial/knn/detail/fused_l2_knn.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -446,6 +446,7 @@ __global__ __launch_bounds__(Policy::Nthreads, 2) void fusedL2kNN(const DataT* x
}
}
}
__syncwarp();
const int finalNumVals = raft::shfl(numValsWarpTopK[i], 31);
loadWarpQShmem<Policy, Pair>(heapArr[i], &shDumpKV[0], rowId, numOfNN);
updateSortedWarpQ<Pair, myWarpSelect::kNumWarpQRegisters>(
Expand Down

0 comments on commit 29b3a18

Please sign in to comment.