From 9de7f710e67025cce0dc2e4709500092ea4aa6a8 Mon Sep 17 00:00:00 2001 From: op-hunter Date: Sat, 30 May 2020 12:27:59 +0800 Subject: [PATCH] optimize the implemention of k-selection on faiss gpu version (#2467) * optimize the implemention of k-selection on faiss gpu version Signed-off-by: cmli * update change log, re-classified the pr from feature to improvement Signed-off-by: cmli * update change log Signed-off-by: cmli Co-authored-by: cmli --- CHANGELOG.md | 1 + core/src/index/thirdparty/faiss/gpu/utils/Select.cuh | 12 ++---------- 2 files changed, 3 insertions(+), 10 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 37b5af7ba3..29b2ac9785 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -12,6 +12,7 @@ Please mark all change in change log and use the issue from GitHub ## Improvement - \#2381 Upgrade FAISS to 1.6.3 - \#2441 Improve Knowhere code coverage +- \#2466 optimize k-selection implemention of faiss gpu version ## Task diff --git a/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh b/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh index 43a1cc1893..0dad487140 100644 --- a/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh +++ b/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh @@ -124,16 +124,8 @@ struct BlockSelect { __device__ inline void addThreadQ(K k, V v) { if (Dir ? Comp::gt(k, warpKTop) : Comp::lt(k, warpKTop)) { - // Rotate right -#pragma unroll - for (int i = NumThreadQ - 1; i > 0; --i) { - threadK[i] = threadK[i - 1]; - threadV[i] = threadV[i - 1]; - } - - threadK[0] = k; - threadV[0] = v; - ++numVals; + threadK[numVals] = k; + threadV[numVals ++] = v; } }