milvus-io
/
milvus
mirror of https://github.com/milvus-io/milvus.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Improve IVF query on GPU when no entity deleted (#5204) 

Signed-off-by: shengjun.li <shengjun.li@zilliz.com>
pull/5251/head
shengjun.li 2021-05-14 11:29:57 +08:00 committed by GitHub
parent 2e4aaff531
commit 623ae55ff5
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 26 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

42
core/src/index/thirdparty/faiss/gpu/impl/IVFUtilsSelect1.cu vendored
View File

@ -68,29 +68,39 @@ pass1SelectLists(void** listIndices,
// BlockSelect add cannot be used in a warp divergent circumstance; we
// handle the remainder warp below
for (; i < limit; i += blockDim.x) {
index = getListIndex(queryId,
start + i,
listIndices,
prefixSumOffsets,
topQueryToCentroid,
opt);
if (bitsetEmpty || (!(bitset[index >> 3] & (0x1 << (index & 0x7))))) {
do {
if (!bitsetEmpty) {
index = getListIndex(queryId,
start + i,
listIndices,
prefixSumOffsets,
topQueryToCentroid,
opt);
if (bitset[index >> 3] & (0x1 << (index & 0x7))) {
break;
}
}
heap.addThreadQ(distanceStart[i], start + i);
}
} while(0);
heap.checkThreadQ();
}
// Handle warp divergence separately
if (i < num) {
index = getListIndex(queryId,
start + i,
listIndices,
prefixSumOffsets,
topQueryToCentroid,
opt);
if (bitsetEmpty || (!(bitset[index >> 3] & (0x1 << (index & 0x7))))) {
do {
if (!bitsetEmpty) {
index = getListIndex(queryId,
start + i,
listIndices,
prefixSumOffsets,
topQueryToCentroid,
opt);
if (bitset[index >> 3] & (0x1 << (index & 0x7))) {
break;
}
}
heap.addThreadQ(distanceStart[i], start + i);
}
} while(0);
}
// Merge all final results