From 2d2be1e24262f2a106318a940b60318d464213ad Mon Sep 17 00:00:00 2001 From: zh Wang Date: Wed, 15 Jul 2026 22:00:12 +0800 Subject: [PATCH] Fix missing trailing barrier in ivf_rabitq blockReduceSum --- cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh b/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh index 51d3651b45..6132ee7f55 100644 --- a/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh +++ b/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh @@ -32,6 +32,10 @@ __inline__ __device__ T blockReduceSum(T val) int wid = threadIdx.x >> 5; val = warpReduceSum(val); + // Wait until every warp has finished reading shared[] in the previous call + // before this call overwrites it. Doing this after the warp-local reduction + // overlaps the extra synchronization with useful work. + __syncthreads(); if (lane == 0) shared[wid] = val; __syncthreads();