Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
203 commits
Select commit Hold shift + click to select a range
e90929a
Add library for IVF-RaBitQ
jamxia155 Nov 3, 2025
4a5c6dc
Add benchmarking executables for IVF-RaBitQ
jamxia155 Nov 3, 2025
ec88c25
Add --executable-dir option
tfeher Feb 12, 2025
bcc80e2
Add documentation about how to describe new datasets
tfeher Feb 12, 2025
4db552f
fix style
tfeher Feb 12, 2025
fbfe242
update docstring
tfeher Apr 1, 2025
5c53816
Do not prompt for executable-dir
tfeher Jul 9, 2025
285a9dd
Enable IVF-RaBitQ in cuvs_bench python wrapper
jamxia155 Nov 3, 2025
9f5a3d2
Use SPDX for copyright headers
jamxia155 Nov 3, 2025
133f808
Add documentation for 3rd-party dependency
jamxia155 Nov 3, 2025
09075bc
Add FAISS CPU IVF-RaBitQ algorithm
jamxia155 Nov 4, 2025
4884e89
Enable FAISS CPU IVF-RaBitQ in cuvs_bench python wrapper
jamxia155 Nov 5, 2025
cbd2a05
Rename parameter for consistency
jamxia155 Nov 6, 2025
7249455
Fix cuVS build issues with RaBitQ (#4)
Stardust-SJF Nov 12, 2025
1a26a71
Handle host and device data in build
jamxia155 Nov 12, 2025
d49bd0b
Merge remote-tracking branch 'Stardust-SJF_fork/jamxia_cuvs_ivf_rabit…
jamxia155 Nov 13, 2025
794b421
Disable separable compilation for IVF-RaBitQ code
jamxia155 Nov 13, 2025
f1fc50b
Brev benchmark instructions (#1)
tfeher Nov 13, 2025
928945b
Remove outdated instructions
jamxia155 Nov 13, 2025
9a5c0ef
Plumbing for passing raft handle to IVF-RaBitQ
jamxia155 Nov 18, 2025
389c917
Update rotator_gpu class
jamxia155 Nov 18, 2025
114d560
Migrate RotatorGPU class to RAFT
jamxia155 Nov 19, 2025
db8a437
Remove cuBLAS from RotatorGPU class
jamxia155 Nov 19, 2025
b263628
Remove Eigen dependency in `DataQuantizerGPU`
jamxia155 Nov 20, 2025
8404b99
Remove uses of Eigen library
jamxia155 Nov 20, 2025
035e978
Remove dependency `Eigen`
jamxia155 Nov 20, 2025
c73ef60
(WIP) Add tests for IVF-RaBitQ
jamxia155 Nov 21, 2025
6cac52d
Merge remote-tracking branch 'upstream/main' into jamxia_cuvs_ivf_rabitq
jamxia155 Nov 21, 2025
f03ec47
Replace header guards with `#pragma once`
jamxia155 Nov 21, 2025
f0be854
Add namespace
jamxia155 Nov 22, 2025
1c444c7
Add tests for IVF-RaBitQ
jamxia155 Nov 24, 2025
a680ca7
Check errors on CUDA API calls and kernel launches
jamxia155 Nov 24, 2025
596997d
Rename class member cudaStream_t in RotatorGPU
jamxia155 Nov 24, 2025
1cb8299
Migrate IVFGPU to async CUDA calls and launches
jamxia155 Nov 24, 2025
ef476ba
Avoid using default stream
jamxia155 Nov 24, 2025
a844b9b
Use async CUDA calls in InitializerGPU classes
jamxia155 Nov 25, 2025
aa0a195
Use async CUDA calls for device results pool
jamxia155 Nov 25, 2025
b2f8a7b
Use async calls in DataQuantizerGPU class
jamxia155 Nov 25, 2025
e2d5d1a
Use async CUDA calls in BatchedQueryGatherer class
jamxia155 Nov 25, 2025
c698e18
Use async calls in SearcherGPU class
jamxia155 Nov 25, 2025
35e9249
Add class members for resource handle and stream
jamxia155 Nov 26, 2025
c381ab2
Clean up member ownership and access in IVFGPU
jamxia155 Nov 26, 2025
abc657b
Use RAFT containers in RatatorGPU class
jamxia155 Nov 26, 2025
0883c37
Use RAFT containers in InitializerGPU class
jamxia155 Nov 26, 2025
c7ecd1c
Change mdarray index type to int64_t
jamxia155 Nov 27, 2025
55bb4fa
Use RAFT containers in IVFGPU class
jamxia155 Nov 27, 2025
f2d159a
Use RAFT containers in DeviceResultPool struct
jamxia155 Dec 2, 2025
cc28ed8
Use RAFT containers and smart pointers in SearcherGPU class
jamxia155 Dec 3, 2025
384da65
Move IVF-RaBitQ internal headers to cpp/src
jamxia155 Dec 3, 2025
0885531
Revert "Move IVF-RaBitQ internal headers to cpp/src"
jamxia155 Dec 3, 2025
6350521
Synchronize with the updates of IVF-RaBitQ-GPU. (#6)
Stardust-SJF Dec 3, 2025
0421efa
Move IVF-RaBitQ internal headers to cpp/src
jamxia155 Dec 3, 2025
984f2df
Merge remote-tracking branch 'Stardust-SJF_fork/jamxia_cuvs_ivf_rabit…
jamxia155 Dec 3, 2025
5ac201a
Fix a bug
jamxia155 Dec 3, 2025
b10ecae
Remove debug code
jamxia155 Dec 3, 2025
f0a61d6
Remove commented-out code
jamxia155 Dec 4, 2025
f0dc124
Fix memory leaks
jamxia155 Dec 4, 2025
f587114
Initialize elements in padded queries
jamxia155 Dec 4, 2025
3214bca
Add default initializations for class members
jamxia155 Dec 4, 2025
b4970f7
Remove unused declarations
jamxia155 Dec 4, 2025
d128515
Remove unused utils code
jamxia155 Dec 5, 2025
5bf3215
Only create padded queries matrix if needed
jamxia155 Dec 5, 2025
0548b54
Merge remote-tracking branch 'upstream/main' into jamxia_cuvs_ivf_rabitq
jamxia155 Dec 5, 2025
ef94d1f
Remove unused code in memory.hpp
jamxia155 Dec 5, 2025
14c9df5
Remove unused code in InitializerGPU class
jamxia155 Dec 5, 2025
5fc2e00
Remove BatchedQueryGatherer class
jamxia155 Dec 5, 2025
1df4385
Remove unused code from IVFGPU class
jamxia155 Dec 5, 2025
7bff107
Remove unused code in DataQuantizerGPU class
jamxia155 Dec 5, 2025
9ee467e
Remove unused code in SearcherGPU class
jamxia155 Dec 5, 2025
dcd36d3
Remove pool_gpu.cu/.cuh (no longer used)
jamxia155 Dec 5, 2025
322df78
Update API for 1-bit quantization support
jamxia155 Dec 5, 2025
8a3096e
Replace cuBLAS calls with templated RAFT wrapper
jamxia155 Dec 6, 2025
443f1ca
Set max dynamic shared mem size as needed
jamxia155 Dec 8, 2025
fb839c0
Support 1-bit search for all search modes (#7)
Stardust-SJF Dec 10, 2025
6aee332
Update API for 1-bit RaBitQ
jamxia155 Dec 10, 2025
e401067
Split up SearcherGPU impl
jamxia155 Dec 10, 2025
9402e9e
Consolidate kernel parameters
jamxia155 Dec 10, 2025
0d5d678
Remove a comment
jamxia155 Dec 11, 2025
80b9fac
Enable handling of large top-k value (up to 16384).
jamxia155 Dec 11, 2025
a43fda2
Check that topk value is below max supported.
jamxia155 Dec 16, 2025
35732b2
Remove commented code
jamxia155 Dec 16, 2025
078fb06
Merge remote-tracking branch 'upstream/main' into jamxia_cuvs_ivf_rabitq
jamxia155 Dec 18, 2025
eb1efbb
Updates after merge from upstream
jamxia155 Dec 18, 2025
1298494
Remove unnecessary code
jamxia155 Dec 18, 2025
3b4e71e
Compute max_cluster_length in load_transposed
jamxia155 Dec 18, 2025
a129b78
Add condition to updating threshold
jamxia155 Dec 22, 2025
475179a
Do not use block sort for large top-k
jamxia155 Dec 23, 2025
ccbd0ff
Fix a bug in thresholding
jamxia155 Dec 24, 2025
87809cf
Remove thresholding for some search code paths
jamxia155 Dec 24, 2025
444058b
Reduce atomics and shared mem in LUT32 search mode
jamxia155 Dec 26, 2025
144a9e9
Reduce atomics and shared mem in LUT16 search mode
jamxia155 Dec 27, 2025
e3af96e
Reduce atomics and shared mem in QUANT4/8 search
jamxia155 Dec 29, 2025
6848b3a
Remove unnecessary shared memory variables
jamxia155 Dec 29, 2025
07e7a88
Replace device mem allocation with RAFT containers
jamxia155 Jan 1, 2026
f93550b
Replace some raw allocations with RAFT containers
jamxia155 Jan 4, 2026
ada5c0b
Minimize allocation for intermediate output
jamxia155 Jan 6, 2026
c056510
Merge remote-tracking branch 'upstream/main' into jamxia_cuvs_ivf_rabitq
jamxia155 Jan 12, 2026
3453328
Rename library
jamxia155 Jan 13, 2026
5ae864b
Optimize recall calculation for large k
jamxia155 Jan 20, 2026
1c87662
Add build API accepting host dataset
jamxia155 Jan 21, 2026
c5e97e8
Enable subsampling of raw dataset for clustering
jamxia155 Feb 3, 2026
03e93db
Change default kmeans_trainset_fraction for tests
jamxia155 Feb 4, 2026
70068eb
Optimize recall calculation
jamxia155 Feb 4, 2026
7634ae3
Preallocate hashset if possible
jamxia155 Feb 5, 2026
29ecc55
Add API for returning index length
jamxia155 Feb 11, 2026
f6e2a86
Improve error reporting
jamxia155 Feb 11, 2026
cbb7154
Merge remote-tracking branch 'upstream/main' into jamxia_cuvs_ivf_rabitq
jamxia155 Feb 12, 2026
0f0c444
Update member declaration order
jamxia155 Feb 17, 2026
8ade5e9
Merge remote-tracking branch 'upstream/main' into jamxia_cuvs_ivf_rabitq
jamxia155 Feb 17, 2026
dfc0708
Merge remote-tracking branch 'upstream/main' into jamxia_cuvs_ivf_rabitq
jamxia155 Feb 22, 2026
bcb230d
Fix #includes
jamxia155 Feb 22, 2026
107659d
Use safely_launch_kernel_with_smem_size
jamxia155 Feb 22, 2026
f4dd5a8
Fix linalg/gemm include
tfeher Feb 25, 2026
1553d56
Revert "Update member declaration order"
jamxia155 Feb 25, 2026
d42b61f
Revert "Preallocate hashset if possible"
jamxia155 Feb 25, 2026
2fd96a1
Revert "Optimize recall calculation"
jamxia155 Feb 25, 2026
9503329
Revert "Optimize recall calculation for large k"
jamxia155 Feb 25, 2026
3ccf526
Refactor JIT LTO kernel generation (#1812)
KyleFromNVIDIA Feb 23, 2026
3f8b76c
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Mar 3, 2026
fb32dcb
Remove feature branch-specific content
jamxia155 Mar 3, 2026
443aa93
Restore cuVS Bench CMakeLists
jamxia155 Mar 3, 2026
fb26176
Fix style check failures
jamxia155 Mar 3, 2026
994e951
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Mar 4, 2026
5684f7b
Remove oudated algo name
jamxia155 Mar 18, 2026
bbeb058
Minor updates based on review comments
jamxia155 Mar 18, 2026
75b7b19
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Mar 18, 2026
adad8bc
Update index build parameters
jamxia155 Mar 18, 2026
dd6b4ba
Update index build parameter default
jamxia155 Mar 18, 2026
088665b
Remove outdated parameter
jamxia155 Mar 18, 2026
04ae19c
Remove unnecessary overload
jamxia155 Mar 18, 2026
4496d93
Let `build` API return the built index
jamxia155 Mar 19, 2026
bef16e9
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Mar 19, 2026
485eade
Revert change to unrelated file
jamxia155 Mar 19, 2026
9761933
Use public header for kmeans clustering
jamxia155 Mar 19, 2026
09c42a5
Remove unnecessary casting
jamxia155 Mar 19, 2026
8ec06e3
Add streaming construction of IVF-RaBitQ index
jamxia155 Mar 21, 2026
f00624b
Add force_streaming parameter IVF-RaBitQ build
jamxia155 Mar 21, 2026
e8f8984
Expose force_streaming in benchmark configuration parser
jamxia155 Mar 21, 2026
9e76b26
Clarify force_streaming only applies to datasets on host
jamxia155 Mar 21, 2026
3543caa
Add forced streaming test for IVF-RaBitQ
jamxia155 Mar 22, 2026
24e1e3d
Code cleanup
jamxia155 Mar 23, 2026
e3ef24b
Refactor IVF-RaBitQ: Remove batch_flag and improve encapsulation
jamxia155 Mar 23, 2026
52793fc
Consolidate quantizer_gpu implementation into single file
jamxia155 Mar 23, 2026
69299d1
Replace CUDA memory calls with RMM/RAFT primitives
jamxia155 Apr 23, 2026
96d9460
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Apr 23, 2026
e28491b
Fixes after merging in main branch
jamxia155 Apr 23, 2026
519e738
Address PR comments
jamxia155 May 5, 2026
4ddf3f9
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 May 5, 2026
b3dbd98
Address review comments
jamxia155 May 5, 2026
858b0be
Refactor duplicated kernel code into computeInnerProductsWithBitwiseB…
jamxia155 May 5, 2026
a33e388
Refactor duplicated kernel code into computeInnerProductsWithBitwise
jamxia155 May 5, 2026
194dae9
Refactor duplicated kernel code into computeInnerProductsWithLUTBlock…
jamxia155 May 5, 2026
c98ec10
Refactor duplicate kernel code into computeInnerProductsWithLUT
jamxia155 May 5, 2026
3576c8f
Refactor duplicate kernel code into computeInnerProductsWithLUT
jamxia155 May 6, 2026
3626313
Refactor duplicate kernel code into computeInnerProductsWithLUT16OptB…
jamxia155 May 6, 2026
b6d64f2
Address review comments
jamxia155 May 6, 2026
7ac97f6
Address review comments
jamxia155 May 6, 2026
02ef1fa
Address review comments
jamxia155 May 6, 2026
a9fc3b9
Address review comments
jamxia155 May 7, 2026
b9e779d
Add error checking for cub calls
jamxia155 May 7, 2026
e8b6b09
Fix cross-stream race in IVF-RaBitQ search path
jamxia155 May 10, 2026
3adea83
Consistently use the same memory resource
jamxia155 May 19, 2026
be19d67
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 May 19, 2026
df876f6
Switch to cuvs::selection::select_k
jamxia155 May 20, 2026
9bd834e
Validate inputs from file; remove dead code
jamxia155 May 20, 2026
ed9d4ab
Switch to raft::linalg::norm
jamxia155 May 20, 2026
3de247a
Cache centroid norms calculation before search
jamxia155 May 20, 2026
66e417a
Replace custom kernel with existing functionality
jamxia155 May 20, 2026
7c0c61a
Pass by reference
jamxia155 May 20, 2026
ee92151
Switch to bulk I/O
jamxia155 May 20, 2026
276ad15
Remove dead code
jamxia155 May 20, 2026
a4385ef
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 May 27, 2026
4582378
Add symbol visibility
jamxia155 May 28, 2026
1ac6136
Lift searcher_gpu_common.cuh symbols out of anonymous namespace
jamxia155 May 28, 2026
12fbcd0
refactor(ivf_rabitq): add JIT-LTO A/B dispatch infrastructure
jamxia155 May 29, 2026
e4adf25
refactor(ivf_rabitq): pilot JIT-LTO conversion of computeInnerProduct…
jamxia155 May 30, 2026
59e0cc8
refactor(ivf_rabitq): extract extract_code as shared JIT-LTO fragment
jamxia155 May 30, 2026
c9326bc
refactor(ivf_rabitq): convert LUT16Opt and Bitwise families to JIT-LTO
jamxia155 Jun 1, 2026
c43b464
refactor(ivf_rabitq): merge ivf_rabitq into cuvs_objs
jamxia155 Jun 1, 2026
a1b8804
refactor(ivf_rabitq): extract compute_ip2_from_long_codes_warp as sha…
jamxia155 Jun 2, 2026
6797b72
refactor(ivf_rabitq): extract compute_lut_ip_for_vec as templated JIT…
jamxia155 Jun 2, 2026
cd1f5db
refactor(ivf_rabitq): extract compute_bitwise_1bit_ip_for_vec as shar…
jamxia155 Jun 3, 2026
5a4c7df
refactor(ivf_rabitq): extract update_threshold_atomicmin as shared JI…
jamxia155 Jun 3, 2026
8382c44
refactor(ivf_rabitq): specialize extract_code and compute_ip2_from_lo…
jamxia155 Jun 3, 2026
13d9e41
refactor(ivf_rabitq): extract compute_bitwise_quantized_ip_for_vec as…
jamxia155 Jun 3, 2026
583808c
refactor(ivf_rabitq): align extract_code and compute_ip2_from_long_co…
jamxia155 Jun 3, 2026
8f30881
refactor(ivf_rabitq): drop legacy __global__ kernels and the JIT-LTO …
jamxia155 Jun 3, 2026
2e02b0f
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Jun 4, 2026
f3d8588
Address review comments
jamxia155 Jun 8, 2026
0eac963
Recast device fragments as inline device functions
jamxia155 Jun 8, 2026
327f82c
Address PR comments
jamxia155 Jun 8, 2026
6316b1b
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Jun 8, 2026
351f8ec
Decouple ivf_rabitq block-sort from ivf_flat
jamxia155 Jun 9, 2026
d6e6476
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Jun 9, 2026
ca903cf
Misc. fixes
jamxia155 Jun 10, 2026
26c533c
Disable cuVS IVF-RaBitQ benchmarks for CPU-only build
jamxia155 Jun 10, 2026
a2e0344
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Jun 10, 2026
330b84e
Fix benchmark executable linkage
jamxia155 Jun 10, 2026
87353b1
Turn off device pointer mode for RAFT matmul
jamxia155 Jun 11, 2026
5929527
Fix host-pointer dereference in streaming IVF-RaBitQ build
jamxia155 Jun 11, 2026
d2846cb
Address review comments
jamxia155 Jun 17, 2026
a17cf0d
Merge remote-tracking branch 'upstream/main' into cuvs_ivf_rabitq
jamxia155 Jun 17, 2026
90eaea0
Remove file from old path
jamxia155 Jun 18, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
195 changes: 195 additions & 0 deletions cpp/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -684,6 +684,192 @@ if(NOT BUILD_CPU_ONLY)
OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_pq/increment_score"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
set(ivf_rabitq_ns "cuvs::neighbors::ivf_rabitq::detail")
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_inner_products_with_lut"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_compute_inner_products_with_lut"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_inner_products_with_lut"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_lut_emit_distances_@with_ex_descriptor@"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/lut_emit_distances_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/lut_emit_distances_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_lut_emit_distances<@with_ex_value@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/lut_emit_distances"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_inner_products_with_lut_block_sort"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut_block_sort_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut_block_sort_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_compute_inner_products_with_lut_block_sort"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_inner_products_with_lut_block_sort"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_lut_block_sort_emit_topk_@with_ex_descriptor@"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/lut_block_sort_emit_topk_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/lut_block_sort_emit_topk_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_lut_block_sort_emit_topk<@with_ex_value@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/lut_block_sort_emit_topk"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_extract_code_@ex_bits@bit"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/extract_code_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/extract_code_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_extract_code<@ex_bits@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/extract_code"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_lut_ip_for_vec_@lut_dtype_descriptor@"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_lut_ip_for_vec_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_lut_ip_for_vec_kernel.cu.in"
FRAGMENT_TAG_FORMAT
"${ivf_rabitq_ns}::fragment_tag_compute_lut_ip_for_vec<${ivf_rabitq_ns}::tag_lut_dtype_@lut_dtype_descriptor@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_lut_ip_for_vec"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_bitwise_quantized_ip_for_vec_@num_bits@bit"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_bitwise_quantized_ip_for_vec_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_bitwise_quantized_ip_for_vec_kernel.cu.in"
FRAGMENT_TAG_FORMAT
"${ivf_rabitq_ns}::fragment_tag_compute_bitwise_quantized_ip_for_vec<@num_bits@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_bitwise_quantized_ip_for_vec"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_inner_products_with_lut16_opt"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut16_opt_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut16_opt_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_compute_inner_products_with_lut16_opt"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_inner_products_with_lut16_opt"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_lut16_opt_emit_distances_@with_ex_descriptor@"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/lut16_opt_emit_distances_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/lut16_opt_emit_distances_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_lut16_opt_emit_distances<@with_ex_value@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/lut16_opt_emit_distances"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_inner_products_with_lut16_opt_block_sort_@with_ex_descriptor@"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut16_opt_block_sort_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_lut16_opt_block_sort_kernel.cu.in"
FRAGMENT_TAG_FORMAT
"${ivf_rabitq_ns}::fragment_tag_compute_inner_products_with_lut16_opt_block_sort<@with_ex_value@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_inner_products_with_lut16_opt_block_sort"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_inner_products_with_bitwise"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_bitwise_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_bitwise_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_compute_inner_products_with_bitwise"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_inner_products_with_bitwise"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_bitwise_emit_distances_@with_ex_descriptor@"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/bitwise_emit_distances_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/bitwise_emit_distances_kernel.cu.in"
FRAGMENT_TAG_FORMAT "${ivf_rabitq_ns}::fragment_tag_bitwise_emit_distances<@with_ex_value@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/bitwise_emit_distances"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_compute_inner_products_with_bitwise_block_sort"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_bitwise_block_sort_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/compute_inner_products_with_bitwise_block_sort_kernel.cu.in"
FRAGMENT_TAG_FORMAT
"${ivf_rabitq_ns}::fragment_tag_compute_inner_products_with_bitwise_block_sort"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/compute_inner_products_with_bitwise_block_sort"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
generate_jit_lto_kernels(
jit_lto_files
NAME_FORMAT "ivf_rabitq_bitwise_block_sort_emit_topk_@with_ex_descriptor@"
MATRIX_JSON_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/bitwise_block_sort_emit_topk_matrix.json"
KERNEL_INPUT_FILE
"${CMAKE_CURRENT_SOURCE_DIR}/src/neighbors/ivf_rabitq/jit_lto_kernels/bitwise_block_sort_emit_topk_kernel.cu.in"
FRAGMENT_TAG_FORMAT
"${ivf_rabitq_ns}::fragment_tag_bitwise_block_sort_emit_topk<@with_ex_value@>"
FRAGMENT_TAG_HEADER_FILES "<cuvs/detail/jit_lto/ivf_rabitq/ivf_rabitq_fragments.hpp>"
OUTPUT_DIRECTORY
"${CMAKE_CURRENT_BINARY_DIR}/generated_kernels/ivf_rabitq/bitwise_block_sort_emit_topk"
KERNEL_LINK_LIBRARIES jit_lto_kernel_usage_requirements
)
set(cagra_ns "cuvs::neighbors::cagra::detail")
generate_jit_lto_kernels(
jit_lto_files
Expand Down Expand Up @@ -1200,6 +1386,15 @@ if(NOT BUILD_CPU_ONLY)
src/neighbors/ivf_pq/detail/ivf_pq_process_and_fill_codes.cu
${ivf_pq_search_inst_files}
${ivf_pq_transform_inst_files}
src/neighbors/ivf_rabitq.cu
src/neighbors/ivf_rabitq/gpu_index/ivf_gpu.cu
src/neighbors/ivf_rabitq/gpu_index/initializer_gpu.cu
src/neighbors/ivf_rabitq/gpu_index/quantizer_gpu.cu
src/neighbors/ivf_rabitq/gpu_index/rotator_gpu.cu
src/neighbors/ivf_rabitq/gpu_index/searcher_gpu.cu
src/neighbors/ivf_rabitq/gpu_index/searcher_gpu_shared_mem_opt.cu
src/neighbors/ivf_rabitq/gpu_index/searcher_gpu_quantize_query.cu
src/neighbors/ivf_rabitq/utils/searcher_gpu_utils.cu
src/neighbors/ivf_sq_index.cpp
src/neighbors/ivf_sq/ivf_sq_build_extend_float_uint8_t_int64_t.cu
src/neighbors/ivf_sq/ivf_sq_build_extend_half_uint8_t_int64_t.cu
Expand Down
18 changes: 18 additions & 0 deletions cpp/bench/ann/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,9 @@ option(CUVS_ANN_BENCH_USE_FAISS_CPU_IVF_FLAT "Include faiss' cpu ivf flat algori
ON
)
option(CUVS_ANN_BENCH_USE_FAISS_CPU_IVF_PQ "Include faiss' cpu ivf pq algorithm in benchmark" ON)
option(CUVS_ANN_BENCH_USE_FAISS_CPU_IVF_RABITQ
"Include faiss' cpu ivf rabitq algorithm in benchmark" ON
)
option(CUVS_ANN_BENCH_USE_FAISS_CPU_IVF_SQ "Include faiss' cpu ivf sq algorithm in benchmark" ON)
option(CUVS_ANN_BENCH_USE_FAISS_CPU_HNSW_FLAT "Include faiss' hnsw algorithm in benchmark" ON)
option(CUVS_ANN_BENCH_USE_CUVS_IVF_FLAT "Include cuVS ivf flat algorithm in benchmark" ON)
Expand All @@ -48,6 +51,7 @@ option(CUVS_ANN_BENCH_SINGLE_EXE
"Make a single executable with benchmark as shared library modules" OFF
)
option(CUVS_KNN_BENCH_USE_CUVS_BRUTE_FORCE "Include cuVS brute force knn in benchmark" ON)
option(CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ "Include cuVS ivf RaBitQ algorithm in benchmark" ON)

# ##################################################################################################
# * Process options ----------------------------------------------------------
Expand Down Expand Up @@ -85,6 +89,7 @@ if(BUILD_CPU_ONLY)
set(CUVS_ANN_BENCH_USE_CUVS_IVF_FLAT OFF)
set(CUVS_ANN_BENCH_USE_CUVS_IVF_SQ OFF)
set(CUVS_ANN_BENCH_USE_CUVS_IVF_PQ OFF)
set(CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ OFF)
set(CUVS_ANN_BENCH_USE_CUVS_CAGRA OFF)
set(CUVS_ANN_BENCH_USE_CUVS_BRUTE_FORCE OFF)
set(CUVS_ANN_BENCH_USE_CUVS_CAGRA_HNSWLIB OFF)
Expand All @@ -102,6 +107,7 @@ if(CUVS_ANN_BENCH_USE_CUVS_IVF_PQ
OR CUVS_ANN_BENCH_USE_CUVS_BRUTE_FORCE
OR CUVS_ANN_BENCH_USE_CUVS_IVF_FLAT
OR CUVS_ANN_BENCH_USE_CUVS_IVF_SQ
OR CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ
OR CUVS_ANN_BENCH_USE_CUVS_CAGRA
OR CUVS_ANN_BENCH_USE_CUVS_CAGRA_HNSWLIB
OR CUVS_KNN_BENCH_USE_CUVS_BRUTE_FORCE
Expand Down Expand Up @@ -249,6 +255,12 @@ if(CUVS_ANN_BENCH_USE_CUVS_IVF_FLAT)
)
endif()

if(CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ)
ConfigureAnnBench(
NAME CUVS_IVF_RABITQ PATH src/cuvs/cuvs_benchmark.cu src/cuvs/cuvs_ivf_rabitq.cu LINKS cuvs
)
endif()

if(CUVS_ANN_BENCH_USE_CUVS_IVF_SQ)
ConfigureAnnBench(
NAME CUVS_IVF_SQ PATH src/cuvs/cuvs_benchmark.cu src/cuvs/cuvs_ivf_sq.cu LINKS cuvs
Expand Down Expand Up @@ -320,6 +332,12 @@ if(CUVS_ANN_BENCH_USE_FAISS_CPU_IVF_PQ)
)
endif()

if(CUVS_ANN_BENCH_USE_FAISS_CPU_IVF_RABITQ)
ConfigureAnnBench(
NAME FAISS_CPU_IVF_RABITQ PATH src/faiss/faiss_cpu_benchmark.cpp LINKS ${CUVS_FAISS_TARGETS}
)
endif()

if(CUVS_ANN_BENCH_USE_FAISS_CPU_IVF_SQ)
ConfigureAnnBench(
NAME FAISS_CPU_IVF_SQ PATH src/faiss/faiss_cpu_benchmark.cpp LINKS ${CUVS_FAISS_TARGETS}
Expand Down
47 changes: 47 additions & 0 deletions cpp/bench/ann/src/cuvs/cuvs_ann_bench_param_parser.h
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,12 @@ extern template class cuvs::bench::cuvs_ivf_pq<float, int64_t>;
extern template class cuvs::bench::cuvs_ivf_pq<uint8_t, int64_t>;
extern template class cuvs::bench::cuvs_ivf_pq<int8_t, int64_t>;
#endif
#if defined(CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ)
#include "cuvs_ivf_rabitq_wrapper.h"
#endif
#ifdef CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ
extern template class cuvs::bench::cuvs_ivf_rabitq<float, int64_t>;
#endif
#if defined(CUVS_ANN_BENCH_USE_CUVS_CAGRA) || defined(CUVS_ANN_BENCH_USE_CUVS_CAGRA_HNSWLIB) || \
defined(CUVS_ANN_BENCH_USE_CUVS_CAGRA_DISKANN)
#include "cuvs_cagra_wrapper.h"
Expand Down Expand Up @@ -207,6 +213,47 @@ void parse_search_param(const nlohmann::json& conf,
}
#endif

#if defined(CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ)
template <typename T, typename IdxT>
void parse_build_param(const nlohmann::json& conf,
typename cuvs::bench::cuvs_ivf_rabitq<T, IdxT>::build_param& param)
{
if (conf.contains("nlist")) { param.n_lists = conf.at("nlist"); }
if (conf.contains("niter")) { param.kmeans_n_iters = conf.at("niter"); }
if (conf.contains("max_points_per_cluster")) {
param.max_train_points_per_cluster = conf.at("max_points_per_cluster");
}
if (conf.contains("bits_per_dim")) { param.bits_per_dim = conf.at("bits_per_dim"); }
if (conf.contains("fast_quantize_flag")) {
param.fast_quantize_flag = conf.at("fast_quantize_flag");
}
if (conf.contains("force_streaming")) { param.force_streaming = conf.at("force_streaming"); }
}

template <typename T, typename IdxT>
void parse_search_param(const nlohmann::json& conf,
typename cuvs::bench::cuvs_ivf_rabitq<T, IdxT>::search_param& param)
{
if (conf.contains("nprobe")) { param.rabitq_param.n_probes = conf.at("nprobe"); }

if (conf.contains("mode")) {
std::string mode = conf.at("mode");
if (mode == "lut16") {
param.rabitq_param.mode = cuvs::neighbors::ivf_rabitq::search_mode::LUT16;
} else if (mode == "lut32") {
param.rabitq_param.mode = cuvs::neighbors::ivf_rabitq::search_mode::LUT32;
} else if (mode == "quant4") {
param.rabitq_param.mode = cuvs::neighbors::ivf_rabitq::search_mode::QUANT4;
} else if (mode == "quant8") {
param.rabitq_param.mode = cuvs::neighbors::ivf_rabitq::search_mode::QUANT8;
} else {
throw std::runtime_error("mode: '" + mode +
"', should be either 'lut16', 'lut32', 'quant4' or 'quant8'");
}
}
}
#endif

#if defined(CUVS_ANN_BENCH_USE_CUVS_CAGRA) || defined(CUVS_ANN_BENCH_USE_CUVS_CAGRA_HNSWLIB) || \
defined(CUVS_ANN_BENCH_USE_CUVS_MG) || defined(CUVS_ANN_BENCH_USE_CUVS_CAGRA_DISKANN)
template <typename T, typename IdxT>
Expand Down
19 changes: 19 additions & 0 deletions cpp/bench/ann/src/cuvs/cuvs_benchmark.cu
Original file line number Diff line number Diff line change
Expand Up @@ -100,6 +100,15 @@ auto create_algo(const std::string& algo_name,
a = std::make_unique<cuvs::bench::cuvs_ivf_pq<T, int64_t>>(metric, dim, param);
}
#endif
#ifdef CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ
if constexpr (std::is_same_v<T, float>) {
if (algo_name == "cuvs_ivf_rabitq") {
typename cuvs::bench::cuvs_ivf_rabitq<T, int64_t>::build_param param;
parse_build_param<T, int64_t>(conf, param);
a = std::make_unique<cuvs::bench::cuvs_ivf_rabitq<T, int64_t>>(metric, dim, param);
}
}
#endif
#ifdef CUVS_ANN_BENCH_USE_CUVS_CAGRA
if (algo_name == "raft_cagra" || algo_name == "cuvs_cagra") {
typename cuvs::bench::cuvs_cagra<T, uint32_t>::build_param param;
Expand Down Expand Up @@ -176,6 +185,16 @@ auto create_search_param(const std::string& algo_name, const nlohmann::json& con
return param;
}
#endif
#ifdef CUVS_ANN_BENCH_USE_CUVS_IVF_RABITQ
if constexpr (std::is_same_v<T, float>) {
if (algo_name == "cuvs_ivf_rabitq") {
auto param =
std::make_unique<typename cuvs::bench::cuvs_ivf_rabitq<T, int64_t>::search_param>();
parse_search_param<T, int64_t>(conf, *param);
return param;
}
}
#endif
#ifdef CUVS_ANN_BENCH_USE_CUVS_CAGRA
if (algo_name == "raft_cagra" || algo_name == "cuvs_cagra") {
auto param = std::make_unique<typename cuvs::bench::cuvs_cagra<T, uint32_t>::search_param>();
Expand Down
9 changes: 9 additions & 0 deletions cpp/bench/ann/src/cuvs/cuvs_ivf_rabitq.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,9 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
* SPDX-License-Identifier: Apache-2.0
*/
#include "cuvs_ivf_rabitq_wrapper.h"

namespace cuvs::bench {
template class cuvs_ivf_rabitq<float, int64_t>;
} // namespace cuvs::bench
Loading
Loading