[BUG]: static_set with a signed extent reads OOB when hash is 0x80000000

[kevkrist]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug (https://github.com/NVIDIA/cuCollections/issues)

Type of Bug

Runtime Error

Describe the bug

When cuco::static_set uses a signed extent type such as cuco::extent<int32_t>, a valid hash value of 0x80000000u can produce a negative initial probe position.

The affected path narrows the hash to the signed extent type and applies cuda::std::abs. The narrowed value is INT_MIN, whose positive magnitude is not representable by int32_t; the signed absolute-value operation therefore has undefined behavior. In the tested CUDA build, the observable result remains INT_MIN. linear_probing then applies signed remainder, producing a negative slot index.

The public static_set::insert call in the reproducer below reads 32 bytes before its slots allocation. It does not depend on cuDF, RMM, or an application framework.

Reproduced with cuCollections revision
f517bbb1277753b1852dfd388993383e401eaa38.
The two functions involved are unchanged on dev as of this report.

How to Reproduce

Reproducer

1. Save this as repro.cu in a checkout of that revision:

// A valid 0x80000000 hash makes static_set probe before its slots allocation
// when the extent type is int32_t.

#include <cuco/static_set.cuh>
#include <cuda/std/functional>
#include <cuda_runtime_api.h>

#include <cstdint>
#include <cstdio>

#define CUDA_TRY(expr)                                                            \
  do {                                                                            \
    auto const status = (expr);                                                   \
    if (status != cudaSuccess) {                                                  \
      std::fprintf(stderr, "%s failed: %s\n", #expr, cudaGetErrorString(status)); \
      return 1;                                                                   \
    }                                                                             \
  } while (false)

struct high_bit_hash {
  std::uint32_t const* value;

  __device__ std::uint32_t operator()(std::int32_t) const noexcept
  {
    // A device load prevents compile-time folding of the INT_MIN case.
    return *value;
  }
};

int main()
{
  using key_type    = std::int32_t;
  using extent_type = cuco::extent<std::int32_t>;
  using probe_type  = cuco::linear_probing<1, high_bit_hash>;
  using set_type    = cuco::static_set<key_type,
                                       extent_type,
                                       cuda::thread_scope_device,
                                       cuda::std::equal_to<key_type>,
                                       probe_type>;

  constexpr std::uint32_t raw_hash = 0x80000000u;
  constexpr key_type key           = 7;

  std::uint32_t* device_hash{};
  key_type* device_key{};
  CUDA_TRY(cudaMalloc(reinterpret_cast<void**>(&device_hash), sizeof(raw_hash)));
  CUDA_TRY(cudaMalloc(reinterpret_cast<void**>(&device_key), sizeof(key)));
  CUDA_TRY(cudaMemcpy(device_hash, &raw_hash, sizeof(raw_hash), cudaMemcpyHostToDevice));
  CUDA_TRY(cudaMemcpy(device_key, &key, sizeof(key), cudaMemcpyHostToDevice));

  // A non-power-of-two capacity is important: INT_MIN % 8 happens to be zero.
  set_type set{extent_type{10},
               cuco::empty_key<key_type>{-1},
               cuda::std::equal_to<key_type>{},
               probe_type{high_bit_hash{device_hash}}};
  std::printf("hash=0x%08x capacity=%d\n", raw_hash, static_cast<int>(set.capacity()));
  std::fflush(stdout);

  set.insert(device_key, device_key + 1);
  CUDA_TRY(cudaDeviceSynchronize());
  return 0;
}

2. Build directly against the cuCollections headers:

nvcc -std=c++20 --expt-extended-lambda -lineinfo -arch=sm_100 \
  -I include repro.cu -o repro

Change sm_100 to the architecture of the test GPU if necessary.

3. Run it. The bare run already fails; Compute Sanitizer is only needed to localize the access:

$ ./repro
hash=0x80000000 capacity=10
terminate called after throwing an instance of 'cuco::cuda_error'
  what():  CUDA error at: include/cuco/detail/storage/counter_storage.cuh97: cudaErrorIllegalAddress an illegal memory access was encountered
Aborted (core dumped)

4. Run it under Compute Sanitizer:

compute-sanitizer \
  --tool memcheck \
  --error-exitcode 86 \
  --destroy-on-device-error context \
  --print-limit 5 \
  ./repro

Observed output from the command above:

========= COMPUTE-SANITIZER
hash=0x80000000 capacity=10
========= Invalid __global__ read of size 4 bytes
=========     at cuco::bucket_storage_ref<int, (int)1, cuco::valid_extent<int, (unsigned long)18446744073709551615>>::operator [](int) const+0x3e0 in bucket_storage.inl:76
=========     by thread (0,0,0) in block (0,0,0)
=========     Access to 0x5f3fffffe0 is out of bounds
=========     and is 32 bytes before the nearest allocation at 0x5f40000000 of size 44 bytes
=========         Device Frame: ... open_addressing_ref_impl<...>::insert<int>(T1)+0x330 in open_addressing_ref_impl.cuh:389
=========         Device Frame: ... operator_impl<cuco::op::insert_tag, ...>::insert<int>(T1)+0xa0 in static_set_ref.inl:394
=========         Device Frame: ... cuco::detail::open_addressing_ns::insert_if_n<...>(...)+0xa0 in kernels.cuh:83
...
========= ERROR SUMMARY: 3 errors

Compute Sanitizer exits with status 86. The first error is the invalid slot read; the remaining
two are cudaErrorLaunchFailure fallout after the CUDA context is destroyed.

Root-cause analysis

include/cuco/detail/utils.cuh
narrows the hash before taking its signed absolute value:

template <typename SizeType, typename HashType>
__host__ __device__ constexpr SizeType to_positive(HashType hash)
{
  if constexpr (cuda::std::is_signed_v<SizeType>) {
    return cuda::std::abs(static_cast<SizeType>(hash));
  } else {
    return static_cast<SizeType>(hash);
  }
}

sanitize_hash forwards to to_positive, and
linear_probing::make_iterator
then performs the remainder in the signed extent type:

size_type const init = cuco::detail::sanitize_hash<size_type>(hash_(probe_key)) %
                       (upper_bound / BucketSize) * BucketSize;

Evaluating cuco::detail::sanitize_hash<int32_t>(0x80000000u) in a device kernel on this build returns -2147483648, confirming that abs does not make the value positive.

The reproducer uses only the public static_set::insert API. It does not call cuco::detail::sanitize_hash directly.

The same signed remainder appears in the cooperative-group linear_probing::make_iterator overload and in both double_hashing::make_iterator overloads, so those paths are likely affected as well.

Expected behavior

Every value in a hasher's unsigned result range should produce an in-bounds probe position. The one-key insertion should complete without a Compute Sanitizer error.

Reproduction link

No response

Operating System

Ubuntu 24.04.1

nvidia-smi output

GPU: NVIDIA GB200 NVL4, sm_100

NVCC version

Driver: 580.105.08
CUDA compiler: 13.2.78

[sleeepyjack]

Thanks for opening this issue. I did some investigation and there seem to be a few places where signed extents can break. I'll open a PR soon to fix them.