Add small key type tests by ksapru · Pull Request #812 · NVIDIA/cuCollections

ksapru · 2026-05-01T20:55:59Z

Extends the static_set unit-test type matrices to cover sub-4-byte key types (uint8_t and uint16_t), as requested in the good first issue.

Changes:

Added uint8_t and uint16_t to the TEMPLATE_TEST_CASE_SIG type axes in all 8 templated static_set test files
Added a compile-time num_keys guard (sizeof(Key) == 1 ? 100 : 400) to keep key sequences within each type's representable range, leaving room for the 0xFF / 0xFFFF empty sentinel
Adjusted hardcoded gold_capacity values in unique_sequence_test and retrieve_all_test to be conditional on num_keys
Parameterized expected sum values in for_each_test to match the reduced key count for uint8_t

The underlying implementation (packed_cas via cuda::atomic_ref) already supports 1B and 2B atomic operations — these tests verify that correctness holds end-to-end for small key types.

Addresses #805

copy-pr-bot · 2026-05-01T20:56:02Z

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.

sleeepyjack · 2026-05-01T20:58:48Z

/ok to test 7b0beb1

sleeepyjack · 2026-05-01T21:01:02Z

/ok to test 323adf8

ksapru · 2026-05-02T15:47:48Z

/ok to test 281721b

sleeepyjack · 2026-05-02T22:32:34Z

/ok to test 281721b

sleeepyjack · 2026-05-03T15:58:07Z

/ok to test 61993aa

sleeepyjack · 2026-05-03T21:29:55Z

/ok to test 22ad7ae

sleeepyjack · 2026-05-04T00:01:49Z

/ok to test 6dc9527

sleeepyjack · 2026-05-04T00:47:58Z

CI passes. Nice work! I'll give it a review on Monday or Tuesday depending on my workload. In the meantime, we can prepare follow up PRs: (1) replicate the same test extensions to all data structures, also sparsely testing combinations of differenty sized key/value types (e.g. 1B keys and 4B payloads), (2) go over the docs, readme, and asserts to make sure they reflect the new type constraints.

ksapru · 2026-05-04T12:11:37Z

Thanks @sleeepyjack. Is there an easier way to reproduce CI locally in the future?

PointKernel · 2026-05-04T16:27:22Z

                                                  cuda::proclaim_return_type<Key>(build_fn{}));

  set.insert_async(keys_begin, keys_begin + num_keys);
+  CUCO_CUDA_TRY(cudaDeviceSynchronize());


question: why is this needed?

It shouldn't be required since it is running on the default stream IIUC. It's probably a red herring trying to fix open bug #804.

PointKernel · 2026-05-04T16:28:21Z

+template <>
+struct packed<sizeof(uint8_t)> {
+  using type = uint8_t;  ///< Packed type as `uint8_t` if the size of the object is 1
+};


We pack paired data (e.g., a 32-bit key and 32-bit value) into a single 64-bit unit for CAS, used only in map scenarios. Therefore, the smallest CAS we need is 8 bits (1B), and the smallest packed type is 2B, making the 1B overload unnecessary.

Suggested change

template <>

struct packed<sizeof(uint8_t)> {

using type = uint8_t; ///< Packed type as `uint8_t` if the size of the object is 1

};

PointKernel · 2026-05-04T17:00:19Z

+template <>
+struct packed<sizeof(uint16_t)> {
+  using type = uint16_t;  ///< Packed type as `uint16_t` if the size of the object is 2
+};


could we add a 2B key-value pair in the map unit test to exercise this code path? One test case is sufficient.

PointKernel · 2026-05-04T17:10:10Z

Thanks @sleeepyjack. Is there an easier way to reproduce CI locally in the future?

@ksapru All the CI environment is shared via devcontainers. You could use them to repro the issue.

sleeepyjack · 2026-05-05T00:00:40Z

Thanks @sleeepyjack. Is there an easier way to reproduce CI locally in the future?

@ksapru , yes, @PointKernel 's suggestion is the right way to go. Note that if CI fails, the associated test/build log will tell you at the end how to spin up the corresponding devcontainer environment to reproduce the bug. If you don't have access to a GPU system, you can still use the base images of the devcontainers (e.g., this image) directly to at least compile the code locally by running the ci/build.sh script.

sleeepyjack · 2026-05-05T00:01:58Z

                                                  cuda::proclaim_return_type<Key>(build_fn{}));

  set.insert_async(keys_begin, keys_begin + num_keys);
+  CUCO_CUDA_TRY(cudaDeviceSynchronize());


It shouldn't be required since it is running on the default stream IIUC. It's probably a red herring trying to fix open bug #804.

sleeepyjack · 2026-05-05T00:03:31Z

+  (uint8_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint8_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 1),
+  (uint16_t, cuco::test::probe_sequence::double_hashing, 2),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 1),
+  (uint16_t, cuco::test::probe_sequence::linear_probing, 2),


I wonder if using singed integers (as in the rest of the tests) would make sense for consistency. Plus, it would probably also catch some unintended overflow/wrap-around errors. Just a suggestion though - no hard requirement.

ksapru added 2 commits April 29, 2026 12:28

Fix dangling allocator reference in storage deleters

e7f0ce2

test(static_set): add uint8_t and uint16_t to type axes

3d06afd

ksapru marked this pull request as ready for review May 1, 2026 20:56

ksapru requested review from PointKernel and sleeepyjack as code owners May 1, 2026 20:56

NVIDIA deleted a comment from copy-pr-bot Bot May 1, 2026

test(static_set): add additional edge case

e77f232

ksapru force-pushed the add-small-key-type-tests branch from 323adf8 to e77f232 Compare May 2, 2026 14:33

ksapru force-pushed the add-small-key-type-tests branch from 281721b to 390844b Compare May 3, 2026 12:45

Fix sign conversion error

22ad7ae

ksapru force-pushed the add-small-key-type-tests branch from 61993aa to 22ad7ae Compare May 3, 2026 17:01

ksapru and others added 2 commits May 3, 2026 17:48

Merge branch 'dev' into add-small-key-type-tests

72fb2dd

[pre-commit.ci] auto code formatting

6dc9527

PointKernel reviewed May 4, 2026

View reviewed changes

PointKernel added the type: feature request New feature request label May 4, 2026

sleeepyjack reviewed May 5, 2026

View reviewed changes

Conversation

ksapru commented May 1, 2026

Uh oh!

copy-pr-bot Bot commented May 1, 2026

Uh oh!

sleeepyjack commented May 1, 2026

Uh oh!

sleeepyjack commented May 1, 2026

Uh oh!

ksapru commented May 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sleeepyjack commented May 2, 2026

Uh oh!

sleeepyjack commented May 3, 2026

Uh oh!

sleeepyjack commented May 3, 2026

Uh oh!

sleeepyjack commented May 4, 2026

Uh oh!

sleeepyjack commented May 4, 2026

Uh oh!

ksapru commented May 4, 2026

Uh oh!

PointKernel May 4, 2026

Choose a reason for hiding this comment

Uh oh!

sleeepyjack May 5, 2026 • 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!

PointKernel May 4, 2026

Choose a reason for hiding this comment

Uh oh!

PointKernel May 4, 2026

Choose a reason for hiding this comment

Uh oh!

PointKernel commented May 4, 2026

Uh oh!

sleeepyjack commented May 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sleeepyjack May 5, 2026 • 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!

sleeepyjack May 5, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

ksapru commented May 2, 2026 •

edited

Loading

sleeepyjack May 5, 2026 •

edited

Loading

sleeepyjack commented May 5, 2026 •

edited

Loading

sleeepyjack May 5, 2026 •

edited

Loading