From 2d2be1e24262f2a106318a940b60318d464213ad Mon Sep 17 00:00:00 2001 From: zh Wang Date: Wed, 15 Jul 2026 22:00:12 +0800 Subject: [PATCH 1/2] 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(); From 9ba98fd996162de0962c8f36f38b7820eb9ab61e Mon Sep 17 00:00:00 2001 From: aamijar Date: Thu, 23 Jul 2026 04:56:58 +0000 Subject: [PATCH 2/2] fix copyright --- cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh b/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh index 6132ee7f55..d2bb6641e8 100644 --- a/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh +++ b/cpp/src/neighbors/ivf_rabitq/utils/reductions.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */