From cdd40689f531359c3c86ef8cdf45dba32c963ebe Mon Sep 17 00:00:00 2001
From: op-hunter <ophunter52@gmail.com>
Date: Thu, 15 Oct 2020 19:39:49 +0800
Subject: [PATCH] withdraw pr2467 on branch 0.11.0 (#4002)

Signed-off-by: cmli <chengming.li@zilliz.com>

Co-authored-by: cmli <chengming.li@zilliz.com>
---
 core/src/index/thirdparty/faiss/gpu/utils/Select.cuh | 12 ++++++++++--
 1 file changed, 10 insertions(+), 2 deletions(-)

diff --git a/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh b/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh
index 0dad487140..43a1cc1893 100644
--- a/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh
+++ b/core/src/index/thirdparty/faiss/gpu/utils/Select.cuh
@@ -124,8 +124,16 @@ struct BlockSelect {
 
   __device__ inline void addThreadQ(K k, V v) {
     if (Dir ? Comp::gt(k, warpKTop) : Comp::lt(k, warpKTop)) {
-        threadK[numVals] = k;
-        threadV[numVals ++] = v;
+      // 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;
     }
   }