Support Standalone RaBitQ in cuVS by Stardust-SJF · Pull Request #2377 · NVIDIA/cuvs

Stardust-SJF · 2026-07-29T10:31:29Z

This PR introduces the standalone RaBitQ quantization components into cuVS, which provide a convenient way to quantize vectors into RaBitQ codes.

Note that it depends on #2362. Only the commits on top of cuvs_ivf_rabitq_syn_fix_fused are intended to be reviewed in this PR.

Imports the standalone RaBitQ GPU quantizer as a first-class preprocessing component, alongside the existing scalar, binary and pq quantizers. It is independent of cuvs::neighbors::ivf_rabitq and shares no code with it. Public API (cuvs::preprocessing::quantize::rabitq, header-installed with the rest of include/cuvs): rotator_kind { matmul, fht_kac }, delta_kind, params rotator - kind + padded_dim + the rotation state (raft containers) quantizer - params + const_scaling_factor + dim/padded_dim pipeline - composes the two; neither owns the other init_rotator / init_quantizer / init_pipeline rotate, transform, transform_residuals The rotator and the quantizer are independent peer types, mirroring how IVFGPU already holds RotatorGPU and DataQuantizerGPU as siblings; `pipeline` is the pq-like convenience that composes them. The factories take a `dim` rather than a dataset because RaBitQ needs no training: they generate the random rotation and derive the scaling constant from (dim, ex_bits, seed). State is plain caller-owned data in raft containers, following binary::quantizer, so the rotator travels with the codes and no serialize API is required - matching every other preprocessing quantizer, none of which has I/O (persistence lives at the index layer, cf. detail/dataset_serialize.hpp). Import notes: - The imported sources land in detail/rabitq/ essentially verbatim, wrapped in ...::rabitq::detail. Include paths, the enclosing namespace and the error macros are the only edits; kernels and numerics are untouched. - CUDA_CHECK/CUBLAS_CHECK, which printed and called exit(), are now RAFT_CUDA_TRY/RAFT_CUBLAS_TRY so failures throw instead of killing the process. - class RotatorGPU is dissolved: it owned raw cudaMalloc buffers and its own cublasHandle_t. Its state now lives in the public rotator, and detail exposes free launchers over caller-owned buffers, taking the cuBLAS handle and the stream from raft::resources. - The pipeline's scratch buffer comes from the workspace resource instead of cudaMalloc, and its glue kernels run on the resource stream. - The fht_kac segment size is derived from the padded dimension, matching the standalone. Known follow-up: quantize.cu and rescale_search.cu still launch on the default stream. Since cuVS builds with CUDA_API_PER_THREAD_DEFAULT_STREAM that does not implicitly sync with the resource stream, so pipeline.cu bridges them with two TODO(rabitq)-marked host syncs. Threading a stream through those two files removes both. Tested on an RTX PRO 6000 (sm_120a): PREPROCESSING_TEST passes 368/368, including 141 new RaBitQ cases covering both rotator kinds, uint8/uint16 codes, padded/pow2/non-pow2 dims, rotation norm preservation, non-degenerate codes, finite factors and run-to-run determinism.

The imported quantizer entry points launched on the per-thread default stream, which under CUDA_API_PER_THREAD_DEFAULT_STREAM does not implicitly synchronize with the resource stream. That was bridged by four host-side syncs, so every transform() paid two full host<->device round trips and serialized the caller's stream. Thread the stream through instead and delete the bridge. - quantize_{fused,full}_on_residuals and the launch_quantize_* helpers now take cudaStream_t as their first parameter (first, so the existing trailing default arguments still work, and without a default of its own so no caller can fall back to the default stream by accident). Both kernel launches use it. - get_const_scaling_factor takes raft::resources const& first: it needs the stream, the workspace memory resource and the device properties. - Its three cudaMalloc/cudaFree pairs become rmm::device_uvector on the workspace resource, so the scratch comes from the caller's memory resource and is freed on the stream rather than at a device-wide synchronization point. - cudaGetDeviceProperties(&prop, 0) becomes raft::resource::get_device_properties(res). Reading device 0's sharedMemPerBlock to pick block_size was wrong whenever the current device was not device 0, and it was the one CUDA call in the file with no RAFT_CUDA_TRY; the resource accessor is also cached, so this drops a driver query. - Both cub::DeviceReduce::Sum calls take the stream and are wrapped in RAFT_CUDA_TRY. - sync_before_quantize / sync_after_quantize are removed from both pipeline.cu and rabitq.cu along with all four call sites. The transform() path now issues no host synchronization at all. The two remaining syncs are init-only and necessary: the async read-back of the scaling factor in get_const_scaling_factor (it returns a host float) and the state uploads in init_state_matmul / init_state_fht_kac (they copy from host vectors that die at scope exit). No kernel body, no numerics and no rotator behaviour changed. PREPROCESSING_TEST still passes 368/368 on an RTX PRO 6000 (sm_120a), unchanged from before.

Aligns the public factories with the verb cuVS already uses for functions that return a new object, and with the split ivf_rabitq itself observes: - make_* returns a new object. ivf_rabitq's internals use it for exactly that (make_compute_inner_products_with_lut_launcher and friends), and cuVS's public headers use it ~106 times across make_const_mdspan, make_scalar_view, make_list_extents, make_aligned_dataset, make_strided_dataset and others. - init_* populates existing state and returns void, e.g. void IVFGPU::init_clusters(cluster_sizes). Every other init_* in the tree is a detail-level seeder (init_centroids, init_plus_plus, init_random_graph, init_key), never public API. Before this change init_rotator / init_quantizer / init_pipeline were the only init_* functions in the whole public API surface, and they return objects, so make_* is both the conventional and the more accurate verb. The detail layer keeps init_state_matmul / init_state_fht_kac: they fill a caller-provided device buffer and return void, which is the init_* semantic. PREPROCESSING_TEST still passes 368/368; no behaviour change.

Brings rabitq to parity with scalar/binary/pq, which all expose a C API under c/. 20 functions, each returning cuvsError_t with its body wrapped in cuvs::core::translate_exceptions: - cuvsRabitqQuantizerParams + ParamsCreate/ParamsDestroy, defaults mirroring the C++ member initializers. - Opaque cuvsRabitqRotator / cuvsRabitqQuantizer handles + Create/Destroy. - cuvsRabitqRotatorMake / QuantizerMake / PipelineMake. - cuvsRabitqRotate, cuvsRabitqTransform[Full], cuvsRabitqTransformResiduals[Full]. - Getters: rotator kind/padded_dim/state, quantizer dim/padded_dim/ const_scaling_factor. Notes on the shape: - No opaque pipeline handle. The C++ pipeline only composes the two peers, so cuvsRabitqPipelineMake takes two out-parameters instead of adding a third handle type to unwrap. - The handles carry only { uintptr_t addr }. pq and binary also carry a DLDataType because their C++ class is quantizer<T> and the field records which instantiation to cast back to; rabitq's rotator and quantizer are non-templated, so the field would be permanently dead. - The rotator getters are load-bearing rather than conveniences: there is no serialize API by design, so GetKind/GetPaddedDim/GetState are the only way a C caller can keep the rotator alongside the codes it produced. GetState follows cuvsProductQuantizerGetPqCodebook's ownership semantics (a non-owning view of rotator-owned device memory) but its rank and dtype depend on the kind: 2-D float32 for matmul, 1-D uint8 for fht_kac. - Device memory only. pq/scalar/binary dispatch between host and device overloads; every rabitq entry point takes device views, so host tensors are rejected rather than routed. - transform and transform_residuals are each overloaded on their output type alone, which C cannot express, hence the explicit ...Full spellings. - rabitq.h had to be added to c/include/cuvs/core/all.h: header_check.cmake globs c/include and fails an existing test for any header missing from it. Tested on an RTX PRO 6000 (sm_120a): RABITQ_C_TEST passes 2/2, covering both rotator kinds end to end (PipelineMake -> Transform -> GetState -> Destroy), and PREPROCESSING_TEST still passes 368/368.

copy-pr-bot · 2026-07-29T10:31:33Z

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

tfeher

Thank you Jifan for the PR! We are excited to have a standalone RaBitQ quantizer in cuVS. The code is well structured and cleanly implemented. In general, we should aim for more code reuse. Here is my first batch of comments.

tfeher · 2026-10-18T23:18:51Z

+ * @brief RaBitQ quantizer parameters.
+ *
+ * These parameters fully describe the quantizer: RaBitQ is data-oblivious, so
+ * none of them are learned from a dataset (see make_quantizer / make_rotator).


This means there is no clustering involved. We also don't subtract a global mean value from the vectors, right? Do we have assumption about data distribution (e.g. zero mean), or it is not required?

For RaBitQ, it does not need any type of assumption for data distribution, and we also don't subtract a global mean value from the vectors.

Still, we provide a centroid option (something like the global mean) for subtraction. This can provide a tighter error bound in some cases (e.g., for IVF-RaBitQ).

tfeher · 2026-10-18T23:42:10Z

+struct pipeline {
+  /** Rotation applied before quantization. */
+  rabitq::rotator rotator;
+  /** Quantizer applied to the rotated residuals. */
+  rabitq::quantizer quantizer;


Here pipeline is the high level object, the entry point that the user should use for RaBitQ quantizing data. In itself it is logical. But it does not match the naming conventions used in the other quantizers we have. E.g. PQ quantizer is actually a pipeline of multiple steps, still the high level object is called quantizer. I recommend to rename pipeline as "quantizer" and find a different name for the existing quantizer. Tagging @achirkin for suggestions.

Thanks for the suggestion. I've made mainly the following changes:
pipeline → quantizer (the entry point), and the old quantizer → encoder.

tfeher · 2026-10-18T23:49:02Z

+ * A temporary `n_rows x padded_dim` float buffer holding the rotated residuals
+ * is allocated from the workspace resource, so large datasets should be
+ * transformed in batches (a pool workspace resource is recommended).


Most of cuVS algorithms would do this batching under the hood automatically. I think it would make sense to implement it here too.

The algorithms currently only accept device input data. Large host dataset would need manual batching and copying. We should consider accepting host input data as well.

Thanks for the comments. I made the changes so that the quantizer supports both internal batching and host/device input data.

tfeher · 2026-10-19T00:11:47Z

+// Pipeline: pad(dim -> rotator.padded_dim) -> rotate -> [subtract rotated
+// centroid] -> fused quantization on the rotated residuals.
+//
+// Design note: the centroid is rotated with a separate 1-row rotate() call


Won't this make kernel launch overhead a problem?

I think this will not be a problem as launches are lightweight (μs) and only matter for small batches, whereas in most cases, the bottleneck should be the bandwidth for the following subtraction.

By the way, it can be possible to optimize this pipeline, such as subtracting before the rotation, or using a fused kernel.

tfeher · 2026-10-19T00:26:05Z

+// conventions. Numerically stable enough for the few-thousand dimensions used
+// by RaBitQ. Output is row-major.
+// ---------------------------------------------------------------------------
+std::vector<float> random_orthogonal_matrix(int64_t D, uint64_t seed)


I believe ivf_pq::hepers::make_rotation_matrix could do the same task. We should use that instead. (Ideally we should move the helper to preprocessing, and use it both from ivf_pq and ivf_rabitq).

Thanks for the suggestion. I have changed the api to make_rotation_matrix and it turns out to be faster.

tfeher · 2026-10-19T00:40:53Z

+    auto stream = raft::resource::get_cuda_stream(res);
+    constexpr int kBlock = 256;
+    const size_t total = N * D;
+    const int grid = static_cast<int>((total + kBlock - 1) / kBlock);


Suggested change

const int grid = static_cast<int>((total + kBlock - 1) / kBlock);

const int grid = raft::ceildiv<int>(total, kBlock);

Now this handwritten grid disappeared.

tfeher · 2026-10-19T00:53:44Z

+        pad_and_rotate(res, d_centroid, 1, dim, rotator, d_rotated_centroid);
+        constexpr int kBlock = 256;
+        const size_t total = N * D;
+        subtract_row_kernel<<<static_cast<int>((total + kBlock - 1) / kBlock), kBlock, 0, stream>>>(


This can be expressed with raft::linewise_op:

auto rows = raft::make_device_matrix_view<float, int64_t>(d_residuals, N, D); raft::matrix::linewise_op<raft::Apply::ALONG_ROWS>( res, raft::make_const_mdspan(rows), rows, raft::sub_op{}, raft::make_device_vector_view<const float, int64_t>(d_rotated_centroid, dim));

Now expressed with raft::linewise_op:

tfeher · 2026-10-19T00:59:17Z

+
+namespace {
+
+__global__ void pad_rows_kernel(const float* __restrict__ d_src,


This can be expressed with raft::linalg::map:

void pad_rows(raft::resources const& res, const float* d_src, float* d_dst, size_t N, size_t dim, size_t D) { raft::linalg::map_offset( res, raft::make_device_vector_view<float, size_t>(d_dst, N * D), [d_src, dim, D] __device__(size_t idx) { const size_t row = idx / D; const size_t col = idx % D; return col < dim ? d_src[row * dim + col] : 0.0f; }); }

Now replaced with raft::linalg::map

achirkin · 2026-10-28T15:42:46Z

Hi @Stardust-SJF , we have recently made some adjustments to the dataset API in #2858 . Let's have a look how we can adapt the RabitQ quantizer and dataset representation to fit well there? In particular, I opened a PR #2953 to refactor the existing BBQ structures, which derive from RabitQ; as a part of this effort I split out the affine_data (codes+shift+scale) that can be re-used in this RabitQ PR as is.

pipeline.cu hand-rolled two elementwise kernels around the rotator. Both are now raft calls, following review feedback: - pad_rows_kernel -> raft::linalg::map_offset over the padded [N, D] output, in a small pad_rows helper because pad_and_rotate needs it in both the in-place and the out-of-place (matmul) branch. - subtract_row_kernel -> raft::matrix::linewise_op<Apply::ALONG_ROWS> with raft::sub_op, in place on the residuals. The vector is the rotated centroid, so its length is the padded D, not dim. ivf_pq_build.cuh subtracts cluster centers with the same in-place call. - With no hand-written launches left, the grid-size arithmetic and the cudaGetLastError checks go too. The old kernels indexed with int, so N * D >= 2^31 overflowed: at 2.2M rows of dim 1000 (D = 1024) transform() aborted on a CUDA error. The raft calls index with size_t / int64_t and handle that size. Tested on an RTX PRO 6000 (sm_120a): - PREPROCESSING_TEST 368/368, RABITQ_C_TEST 2/2. - Codes, factors and rotator state byte-identical to the previous commit over 96 configurations: dim {33, 100, 200, 960, 1000} x {matmul, fht_kac} x ex_bits {3, 8} x factor_mode {0, 1} x {no centroid, centroid} at 5000 rows, plus 1- and 7-row inputs. Also byte-identical on the last 1000 rows of a 2.0M-row, dim-1000 run, the largest size the old kernels handle. - At 2.2M rows the last 1000 rows match the same rows quantized on their own: exactly for fht_kac, and for matmul up to cuBLAS's shape-dependent sgemm rounding, which the old code shows identically at 2.0M rows. - transform() at 1M rows, dim 1000 is 0.2-0.7% slower than before, most of it in the centroid subtraction.

Following review feedback, the matmul rotator now reuses the random rotation IVF-PQ draws instead of its own host-side Gram-Schmidt. - The body of ivf_pq's internal make_rotation_matrix moves verbatim to cpp/src/preprocessing/detail/rotation_matrix.cuh as cuvs::preprocessing::detail::make_rotation_matrix, with the RNG state as an explicit argument. ivf_pq::helpers::make_rotation_matrix forwards with RngState(7), so IVF-PQ's matrices are unchanged. - init_state_matmul calls it with force_random_rotation and RngState(params.seed). That moves an O(D^3) host step to the GPU (make_rotator: 372 ms -> 3.5 ms at D=1024, 8.5 s -> 15 ms at D=3072), improves orthogonality at large D, and uses the full 64-bit seed (the host mt19937 truncated it to 32 bits). - rotate_matmul applies the stored row-major matrix as y = R x (cublasSgemm OP_T), the orientation IVF-PQ uses. The QR factor is not sign-normalized, so the matrix itself is not exactly Haar; in this orientation that only flips the signs of output coordinates, under which RaBitQ's codes mirror and its factors and estimates stay the same. - raft::random::normal sizes its grid by the SM count, so the same seed gives a different matrix on a different GPU model. The C++ and C docs therefore drop the "store the seed and re-run make_rotator" recipe: callers keep the rotator's buffers alongside the codes. The docs also state the y = R x orientation. fht_kac is untouched. Tested on an RTX PRO 6000 (sm_120a): - PREPROCESSING_TEST 368/368, RABITQ_C_TEST 2/2. - ivf_pq::helpers::make_rotation_matrix output byte-identical to the previous commit for five shapes, including the identity and non-square cases. - make_rotator(matmul) equals the helper's draw for the same seed, including a seed above 2^32. - Negating random rows of R leaves transform()'s factors bit-identical and mirrors the codes, for dims 100/128/1000, ex_bits 1/3/7, both output modes, with and without a centroid. - fht_kac transform output byte-identical to the previous commit over 40 configurations; matmul output differs, as expected with a new matrix.

C callers could save a rotator (GetKind + GetState) but not restore it, while C++ callers restore one by filling the rotator's public fields. Re-running cuvsRabitqRotatorMake with the seed used to cover this, but the matmul rotation now depends on the GPU model, so C needs an explicit restore path to stay on par with C++. cuvsRabitqRotatorMakeFromState(res, kind, state, rotator) fills an empty handle, like cuvsRabitqRotatorMake. It copies `state`, which may live in device or host memory, into the rotator's own buffers, infers padded_dim from its shape and rejects state that does not match `kind` (float32 [D, D] for matmul, uint8 [4 * D / 8] for fht_kac). It synchronizes before returning, so `state` may be freed right away. The handle, Create and GetKind docs now name it as the restore half of GetState. The C API grows from 20 to 21 functions; nothing existing changes. Tested on an RTX PRO 6000 (sm_120a): RABITQ_C_TEST 2/2 with a new round trip for both rotator kinds. Restoring from the device state and from a host copy both rotate bit-identically to the original, and state of the other kind is rejected. A deliberately broken copy (one element short) fails that check on both paths. PREPROCESSING_TEST still passes 368/368.

Following review feedback, transform() no longer quantizes the whole dataset in one pass with an n_rows x padded_dim residual scratch: both overloads now loop over row batches, handing detail::quantize_data / quantize_data_full offset pointers into the caller's input and outputs. The detail pipeline is untouched, and rows are quantized independently, so batching only bounds the scratch. The batch size (transform_batch_rows) is chosen for speed, not only memory. Every stage after the input read (pad, rotate, centroid subtraction, quantize) sweeps the batch's residual scratch, so keeping it L2-resident is faster than one unbatched pass: - fht_kac, and matmul with padded_dim <= 256: the scratch is sized to ~3/8 of the L2 cache, clamped to [2048, 65536] rows ([2048, 32768] for matmul) and rounded down to a power of two (64K rows up to dim 128, 32K at 384, 8K around 1000, 4K at 3072 on a 128 MB L2). - wider matmul is GEMM-bound and runs best on large batches: 131072 rows. - as in IVF-PQ's extend, the batch is then halved (down to 128 rows) until its scratch fits half of the free workspace; the other half covers the rotated centroid, the matmul padding temporary and fragmentation. The constants come from a sweep with one process per measurement; comparing batch sizes inside a single process made the unbatched call look 10-15% slower than it is. transform() no longer needs the n_rows <= INT_MAX check, since every kernel launch now sees at most one batch. The docs drop the advice to batch by hand. Tested on an RTX PRO 6000 (sm_120a, 128 MB L2): - PREPROCESSING_TEST 369/369 (a new test checks that batched fht_kac output is bit-identical to rotate() + transform_residuals() over all rows, across several batches plus a short tail that runs the other FHT kernel, both output modes, with and without a centroid), RABITQ_C_TEST 2/2. - Output byte-identical to the previous commit for 40 configurations at 200000 rows (fht_kac and matmul, dims 100/128/960/1000/3000). - One transform() over a ~2 GiB dataset, 56 configurations, fresh processes x3: fht_kac 1.33x faster (geomean; 1.08-1.68x), matmul 1.05x (0.97-1.29x; dim 384 without a centroid is the one 3% regression). - With a 256 MiB workspace and a ~2 GiB dataset the previous commit runs out of memory; this one completes, for both rotators with and without centroid.

sa_quantize_fused_kernel runs one block per row with an int vec_id and wrote the full-factor outputs at d_factors[vec_id * 3 + k], computed in 32-bit int. For vec_id > 715,827,882 (INT_MAX / 3) that overflows: rows up to ~1.43G write before the factors buffer (an illegal-address error or silent corruption), and rows beyond that wrap around and overwrite the factors of earlier rows. Codes and the (delta, vl) outputs were not affected; the kernel's other per-row offsets were already widened to size_t. transform_residuals(..., factors) and cuvsRabitqTransformResidualsFull allow up to INT_MAX rows per call, so they could reach this (it needs >= ~123 GB of device-accessible memory: an H200/B200 or managed memory). transform() is no longer exposed since it batches rows internally. Widen the index like the other per-row offsets in the kernel. PREPROCESSING_TEST 369/369 and RABITQ_C_TEST 2/2 on an RTX PRO 6000 (sm_120a); transform() output byte-identical to the previous commit for 40 configurations at 200000 rows (fht_kac and matmul, both output modes, with and without a centroid).

Following review feedback, transform() now also takes the dataset in host memory, so a dataset larger than the device no longer has to be batched and copied by hand. This follows pq::transform: four new overloads take a raft::host_matrix_view<const float> dataset, while the centroid and all outputs stay on device and the work stays on the GPU. Host rows are copied to the device by batch_load_iterator, in the same internal batches that device input uses, into workspace buffers that transform_batch_rows now counts in its budget. As in IVF-PQ's extend and pq::transform, the copies run on the main stream, or, when res has a stream pool, on a pool stream into two buffers so that the next batch is copied while the current one is quantized. With prefetch on, res is also synced after each batch (the iterator frees its last buffer on the pool stream, as IVF-SQ's extend handles) and once on entry (the pool stream is not ordered after work already queued on res's stream that may still be writing the host rows). Device input keeps the plain loop: the iterator's destructor syncs even when it does not copy. The C API keeps its functions: cuvsRabitqTransform and cuvsRabitqTransformFull now accept a kDLCPU dataset, checking for device memory first as cuvsProductQuantizerTransform does, so pinned and managed tensors are still read in place. The centroid and outputs must still be on device. A non-const device view now matches both the device and the host overloads, as for pq::transform, so the usage examples pass raft::make_const_mdspan. Tested on an RTX PRO 6000 (sm_120a, PCIe Gen5): - PREPROCESSING_TEST 369/369, RABITQ_C_TEST 2/2. The sweep compares host and device input bit for bit; the batched test runs device input, pageable host input with and without a stream pool, and pinned host input with a stream pool against the unbatched reference, overwriting the host rows with NaN right after each call; the C test covers a kDLCPU dataset for both transforms and rejects host codes. Writing every host batch at row 0 is caught by the batched test. - Device input: output byte-identical to the previous commit for 40 configurations at 200000 rows, and speed unchanged (within 0.3%, 16 configurations, fresh processes x3). - Host input, 4 GiB, 7 configurations: output byte-identical to device input for fht_kac and matmul. With a stream pool, pageable input runs at 23-25 GB/s, 2-6% slower than one plain cudaMemcpy of the dataset (24-26 GB/s), and pinned input at 45-56 GB/s (cudaMemcpy: 57 GB/s), which is 3.3-9.5x the device-input time for fht_kac and 1.2-2.5x for matmul. The stream pool makes host input 1.0-1.3x faster for fht_kac and 1.1-1.6x for matmul.

Following review feedback, the high-level object now follows the naming of the other cuVS quantizers: pipeline becomes quantizer (built by make_quantizer), and the former quantizer, which encodes rotated residuals, becomes encoder (make_encoder), matching RaBitQ-Library's rotate -> encode -> estimate split. The quantizer's member is now quantizer::encoder, and the transform() and transform_residuals() overloads take an encoder. The C API follows: cuvsRabitqQuantizer_t and its Create/Destroy/Make/Get* functions become cuvsRabitqEncoder*, and cuvsRabitqPipelineMake becomes cuvsRabitqQuantizerMake. The parameter structs, the rotator API and the transform function names are unchanged. No behaviour changes. PREPROCESSING_TEST 369/369, RABITQ_C_TEST 2/2.

Stardust-SJF added 4 commits July 26, 2026 20:37

Stardust-SJF requested review from a team as code owners July 29, 2026 10:31

github-project-automation Bot added this to Unstructured Data Processing Jul 29, 2026

Stardust-SJF changed the base branch from main to pull-request/2362 July 29, 2026 11:39

cjnolet added improvement Improves an existing functionality non-breaking Introduces a non-breaking change labels Jul 29, 2026

cjnolet added this to Unstructured Data Processing Jul 29, 2026

cjnolet moved this to In Progress in Unstructured Data Processing Jul 29, 2026

cjnolet assigned Stardust-SJF Jul 29, 2026

tfeher mentioned this pull request Sep 21, 2026

Add support for BBQ compressed datasets in NN-Descent and CAGRA #2654

Merged

This was referenced Oct 10, 2026

[CI] Flaky test sweep (Sep–Oct 2026): root causes and fix PRs #2783

Open

Fix IVF-RaBitQ search for queries on a list centroid #2802

Open

tfeher requested changes Oct 19, 2026

View reviewed changes

Stardust-SJF added 7 commits October 29, 2026 10:49

	const int grid = static_cast<int>((total + kBlock - 1) / kBlock);
	const int grid = raft::ceildiv<int>(total, kBlock);


		namespace {

		__global__ void pad_rows_kernel(const float* __restrict__ d_src,

Conversation

Stardust-SJF commented Jul 29, 2026

Uh oh!

copy-pr-bot Bot commented Jul 29, 2026

Uh oh!

tfeher left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

achirkin commented Oct 28, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants