Fix server crashed during search with index pq (#3954)

* fix pq distances cal bug

Signed-off-by: yangxuan <xuan.yang@zilliz.com>

* fix pq distance calculation outofrange according to faiss suggests

Signed-off-by: yangxuan <xuan.yang@zilliz.com>

* fix change log

Signed-off-by: shengjun.li <shengjun.li@zilliz.com>
Co-authored-by: yangxuan <xuan.yang@zilliz.com>

CHANGELOG.md

@@ -4,13 +4,15 @@ Please mark all change in change log and use the issue from GitHub
 
 # Milvus 0.10.4 (TBD)
 ## Bug
--   \#3903 The search of SQ8H index is slower than the previous version.
+-   \#3626 Fix server crash when searching with IVF_PQ on GPU.
+-   \#3903 The performance of IVF_SQ8H in 0.10.3 is degraded.
 -   \#3906 Change DeleteTask state when it is loaded to avoid server crash.
+-   \#4012 Milvus hangs when continually creating and dropping partitions.
 
 ## Feature
 
 ## Improvement
--   \#3775 Improve search performance in the case that no item deleted
+-   \#3775 Improve search performance in the case that no item deleted.
 
 ## Task
 

core/src/index/thirdparty/faiss/gpu/impl/PQCodeDistances-inl.cuh

@@ -51,11 +51,14 @@ pqCodeDistances(Tensor<float, 2, true> queries,
   auto code = threadIdx.x;
   auto subQuantizer = blockIdx.y;
 
+
   // Each thread will load the pq centroid data for the code that it
   // is processing
+  if(!isLoadingThread) {
 #pragma unroll
-  for (int i = 0; i < DimsPerSubQuantizer; ++i) {
-    subQuantizerData[i] = pqCentroids[subQuantizer][i][code].ldg();
+      for (int i = 0; i < DimsPerSubQuantizer; ++i) {
+        subQuantizerData[i] = pqCentroids[subQuantizer][i][code].ldg();
+      }
   }
 
   // Where we store our query vector
@@ -152,6 +155,8 @@ pqCodeDistances(Tensor<float, 2, true> queries,
           }
         }
       } else {
+
+
         // These are the processing threads
         float dist = 0.0f;