Skip to content
This repository has been archived by the owner on Aug 16, 2023. It is now read-only.

Fix Distance.cu bitset idx match #896

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

Fix Distance.cu bitset idx match #896

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion thirdparty/faiss/faiss/gpu/impl/Distance.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -367,10 +367,13 @@ void runDistance(
k,
streams[curStream]);
} else {
// bitset need to match real idx
auto batchBitsetView = bitset.narrow(
0, int(j / 8), bitset.getSize(0) - int(j / 8));
// Write into the intermediate output
runBlockSelect(
distanceBufView,
bitset,
batchBitsetView,
outDistanceBufColView,
outIndexBufColView,
true,
Expand Down
8 changes: 6 additions & 2 deletions thirdparty/faiss/faiss/gpu/utils/BlockSelectKernel.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -237,7 +237,9 @@ __global__ void blockSelectPair(
bool bitsetEmpty = (bitset.getSize(0) == 0);

for (; i < limit; i += ThreadsPerBlock) {
if (bitsetEmpty || (!(bitset[i >> 3] & (0x1 << (i & 0x7))))) {
// bitset need to match real idx
if (bitsetEmpty ||
(!(bitset[inV[row][i] >> 3] & (0x1 << (inV[row][i] & 0x7))))) {
heap.addThreadQ(*inKStart, *inVStart);
}
heap.checkThreadQ();
Expand All @@ -248,7 +250,9 @@ __global__ void blockSelectPair(

// Handle last remainder fraction of a warp of elements
if (i < inK.getSize(1)) {
if (bitsetEmpty || (!(bitset[i >> 3] & (0x1 << (i & 0x7))))) {
// bitset need to match real idx
if (bitsetEmpty ||
(!(bitset[inV[row][i] >> 3] & (0x1 << (inV[row][i] & 0x7))))) {
heap.addThreadQ(*inKStart, *inVStart);
}
}
Expand Down