Removing filter templates in ivfpq for binary size reduction #1211

jinsolp · 2025-08-04T20:15:37Z

Summary

This PR removed filters from templates in ivfpq. Only applied to compute_similarity_kernel in this PR, but can be applied elsewhere (e.g. ivfflat_interleaved_scan) which will be a follow-up PR to this one.

Perf check

Ran the default setting with bitset flter 20 times to check if this slows anything down.
All numbers are averaged across the 20 runs. Profiing done with nsys.

compute_similarity kernel runtime

Think the extra time is coming from the switch statements inside ivf_to_sample_filter_dev's operator(). But not too much of an overhead.

Binary size reduction

From CI reports: 1107.27 MB ->1040.44 MB

lowener · 2025-08-05T10:19:59Z

cpp/include/cuvs/neighbors/common.hpp

 };

+_RAFT_HOST_DEVICE inline none_sample_filter::none_sample_filter()  = default;
+_RAFT_HOST_DEVICE inline none_sample_filter::~none_sample_filter() = default;


Is there a reason to define them as default here and not in the struct?

From the previous CI run, I could see that the build was failing because it wouldn't let me declare + default with the __device__ annotation. It seems like nvcc ignores __host__ __device__ on default functions when they are defaulted at their first declaration within the struct.

cpp/src/neighbors/sample_filter.cuh

achirkin

Hi @jinsolp, thanks for working on this!
I have a request regarding the synchronization behavior of the filter object. I believe we should strive to avoid the sync at all costs where possible to gradually improve the small-batch and dynamic-batch performance across our algorithms.

cpp/src/neighbors/sample_filter.cuh

jinsolp · 2025-08-07T23:57:32Z

Thanks @achirkin @lowener for your feedbacks! In the binary size sync today, @divyegala suggested using the enum + creating filter inside kernel approach. This is a much less complicated approach, and also shows similar perf to what we already have in this PR (launching separate kernels to track lifetime of kernels). Latest commit introduces this approach.

divyegala · 2025-08-08T00:01:37Z

@jinsolp thanks for the changes! IMO the diff looks much cleaner now.

achirkin · 2025-08-08T07:55:48Z

@jinsolp thanks, would you mind to benchmark this for somewhat larger values of k? let's say 64, 128, 256.
The kind of changes introduced in the kernel seem likely to increase the register pressure, which may lead to performance degradation when we already use a lot of registers - that is with a beefier fused top-k component.

cpp/src/neighbors/ivf_pq/ivf_pq_compute_similarity_impl.cuh

cjnolet · 2025-08-08T20:24:25Z

cpp/src/neighbors/ivf_pq/ivf_pq_compute_similarity.cuh

    const uint32_t* index_list,                                                                \
    float* query_kths,                                                                         \
-    IvfSampleFilterT sample_filter,                                                            \
+    const filtering::base_filter& sample_filter,                                               \


@robertmaynard is this going to impact the update you just made for CUDA 13?

jinsolp · 2025-08-11T19:54:28Z

@achirkin This is the result of using different k's. Some interesting observations

This approach is faster for a larger k (or at least I think we can say that there are no concerning drops in perf)
Kernel suddenly gets faster for k=256 compared to k=128. I'm not sure why this is the case.

achirkin · 2025-08-13T17:00:08Z

cpp/src/neighbors/sample_filter.cuh

+    _RAFT_HOST_DEVICE ivf_filter_dev_args_variant() {}
+    _RAFT_HOST_DEVICE ~ivf_filter_dev_args_variant() {}


Can we get away with no explicitly defined constructor/destructor? I have a feeling this looks more complicated than should :)

The compiler complains if I remove these explicit constructor/destructor

error: the default constructor of "cuvs::neighbors::filtering::ivf_filter_dev::ivf_filter_dev_args_variant" cannot be referenced -- it is a deleted function

Looks like this is because the bitset_filter_args_t is not trivially constructible. Is there a way to workaround this?

What line is this error occurring on?

Happens when I get rid of the explicit constructor/destructor for the union (these two lines)

_RAFT_HOST_DEVICE ivf_filter_dev_args_variant() {} _RAFT_HOST_DEVICE ~ivf_filter_dev_args_variant() {}

Can you report what line the compiler says the error occurs on? Must be somewhere during construction of this struct?

line 159 and 164 of sample_filter.cuh. That would be the ivf_filter_dev constructors.

sample_filter.cuh(159): error: the default constructor of "cuvs::neighbors::filtering::ivf_filter_dev::ivf_filter_dev_args_variant" cannot be referenced -- it is a deleted function { ^ sample_filter.cuh(164): error: the default constructor of "cuvs::neighbors::filtering::ivf_filter_dev::ivf_filter_dev_args_variant" cannot be referenced -- it is a deleted function { ^

Artem's comment below is relevant to this. If you simplify the construction it should be possible.

achirkin · 2025-08-13T17:03:13Z

cpp/src/neighbors/sample_filter.cuh

+
+  _RAFT_HOST_DEVICE ivf_filter_dev(none_filter_args_t args = {}) : tag_(FilterType::None)
+  {
+    new (&args_.none_filter_args) none_filter_args_t(args);


Can't we just initialize using the args_(args) same as you initialize the tag_ above (after you remove explicit constructor/destructor)?

@achirkin @divyegala
Changed to use the initializer list. We still need constructors in the union because bitset_filter_args_t is not trivially constructible and therefore does not have a default constructor. But we don't need the destructor anymore.

cpp/src/neighbors/sample_filter.cuh

divyegala

Looks good to me overall. Can you try to remove all explicit constructors and destructors that are specified as default?

divyegala · 2025-08-13T19:36:44Z

cpp/src/neighbors/sample_filter.cuh

+    _RAFT_HOST_DEVICE ivf_filter_dev_args_variant() {}
+    _RAFT_HOST_DEVICE ~ivf_filter_dev_args_variant() {}


cpp/src/neighbors/sample_filter.cuh

…o bs/filter-ivfpq

lowener · 2025-08-18T11:06:38Z

cpp/src/neighbors/sample_filter.cuh


+struct none_filter_args_t {};
+using bitset_filter_args_t =
+  std::tuple<const int64_t* const*, cuvs::core::bitset_view<uint32_t, int64_t>>;


Can this be tuple<const int64_t* const*, filtering::bitset_filter<uint32_t, int64_t>> to construct the bitset_filter object just once instead of every call to operator()?

Unfortunately, we have problems in each of the following scenarios if we want to do so;
If we were to make the filter once...

outside the compute_similairty_kernel: we have to launch separate kernels for the filter construction, which complicates things

inside the compute_similairty_kernel: we don't want to dynamic dispatch inside the kernel, so it is difficult to share common code.

lowener

LGTM

divyegala

LGTM, thanks!

jinsolp · 2025-08-27T17:00:51Z

/merge

jinsolp added 6 commits July 29, 2025 21:01

remove instantiations

f3a4fa1

revert original filters

8e17b88

rm commented out code

7f556be

avoid device side dynamic dispatch

90ff83a

Merge branch 'rapidsai:branch-25.10' into bs/filter-ivfpq

5a9f8be

change var names

c46d827

jinsolp requested review from a team as code owners August 4, 2025 20:15

github-project-automation bot added this to Vector Search, ML, & Data Mining Release Board Aug 4, 2025

github-project-automation bot moved this to Todo in Vector Search, ML, & Data Mining Release Board Aug 4, 2025

jinsolp self-assigned this Aug 4, 2025

jinsolp moved this from Todo to In Progress in Vector Search, ML, & Data Mining Release Board Aug 4, 2025

github-actions bot added cpp CMake labels Aug 4, 2025

jinsolp added non-breaking Introduces a non-breaking change improvement Improves an existing functionality labels Aug 4, 2025

Merge branch 'branch-25.10' into bs/filter-ivfpq

014f3e7

jinsolp mentioned this pull request Aug 4, 2025

Remove filter templates from ivfpq kernels #1176

Closed

separate default declaration

f4644f5

lowener requested changes Aug 5, 2025

View reviewed changes

jinsolp added 3 commits August 5, 2025 22:18

anonymous namespace for init/destruct kernels

35424fe

address comments

8bd21c4

Merge branch 'branch-25.10' into bs/filter-ivfpq

9eefe78

achirkin requested changes Aug 6, 2025

View reviewed changes

cpp/src/neighbors/sample_filter.cuh Outdated Show resolved Hide resolved

jinsolp added 2 commits August 7, 2025 16:26

create inside kernel with enums

0706391

merge commit resolve conflicts

837bcc0

achirkin reviewed Aug 8, 2025

View reviewed changes

cpp/src/neighbors/ivf_pq/ivf_pq_compute_similarity_impl.cuh Outdated Show resolved Hide resolved

cjnolet added the Waiting for review label Aug 8, 2025

cjnolet reviewed Aug 8, 2025

View reviewed changes

Merge branch 'branch-25.10' into bs/filter-ivfpq

99061e8

jinsolp added 3 commits August 12, 2025 19:05

single struct for args

4f98807

Merge branch 'branch-25.10' into bs/filter-ivfpq

2fec231

change arg order

09068b2

achirkin reviewed Aug 13, 2025

View reviewed changes

divyegala reviewed Aug 13, 2025

View reviewed changes

jinsolp added 7 commits August 13, 2025 20:40

address feedback

ff2452f

remove explicit default

6a945da

rm comment

4a314f3

Merge branch 'branch-25.10' into bs/filter-ivfpq

e1ba613

rm destructor and use init list

f6204f6

Merge branch 'bs/filter-ivfpq' of https://github.com/jinsolp/cuvs int…

01c7c21

…o bs/filter-ivfpq

Merge branch 'branch-25.10' into bs/filter-ivfpq

4c5bbf3

lowener reviewed Aug 18, 2025

View reviewed changes

lowener approved these changes Aug 19, 2025

View reviewed changes

Merge branch 'branch-25.10' into bs/filter-ivfpq

24aa2cd

divyegala approved these changes Aug 25, 2025

View reviewed changes

divyegala and others added 2 commits August 25, 2025 19:44

Merge branch 'branch-25.10' into bs/filter-ivfpq

e3fb085

Merge branch 'branch-25.10' into bs/filter-ivfpq

2c1f267

achirkin approved these changes Aug 27, 2025

View reviewed changes

rapids-bot bot merged commit 4f710ec into rapidsai:branch-25.10 Aug 27, 2025
55 checks passed

github-project-automation bot moved this from In Progress to Done in Vector Search, ML, & Data Mining Release Board Aug 27, 2025

jinsolp deleted the bs/filter-ivfpq branch August 27, 2025 20:15

jameslamb mentioned this pull request Aug 27, 2025

Build and test with CUDA 13.0.0 #1273

Merged

divyegala mentioned this pull request Nov 7, 2025

Convert IVF-PQ filter template to runtime #1518

Closed

jinsolp mentioned this pull request Dec 3, 2025

Reduce Binary Size by using runtime filters for interleaved_scan_kernel #1615

Closed

		_RAFT_HOST_DEVICE ivf_filter_dev_args_variant() {}
		_RAFT_HOST_DEVICE ~ivf_filter_dev_args_variant() {}

Removing filter templates in ivfpq for binary size reduction #1211

Removing filter templates in ivfpq for binary size reduction #1211

Uh oh!

Conversation

jinsolp commented Aug 4, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Perf check

Binary size reduction

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

achirkin left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

jinsolp commented Aug 7, 2025

Uh oh!

divyegala commented Aug 8, 2025

Uh oh!

achirkin commented Aug 8, 2025

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jinsolp commented Aug 11, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jinsolp Aug 13, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

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!

jinsolp Aug 14, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

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!

Uh oh!

divyegala left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

lowener left a comment

Choose a reason for hiding this comment

Uh oh!

divyegala left a comment

Choose a reason for hiding this comment

Uh oh!

jinsolp commented Aug 27, 2025

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

jinsolp commented Aug 4, 2025 •

edited

Loading

jinsolp Aug 13, 2025 •

edited

Loading

jinsolp Aug 14, 2025 •

edited

Loading