[BUG]: DeviceSegmentedSort fails with cudaErrorIllegalAddress when d_begin_offsets[0] != 0

Open

#2,061 opened on Jul 24, 2024 0 comments 0 reactions 1 assignee Claimed by @elstehle View on GitHub

Dominant language: C++
Stars: 2.5k
Forks: 486
Avg merge: 2d 6h
Merged PRs (30d): 295

Description

### Is this a duplicate?

- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Type of Bug

Runtime Error

### Component

CUB

### Describe the bug

The number of items for `cub::DeviceSegmentedSort::SortKeys` is limited to `max(int32_t)` but the function can be called in chunks by manipulating the `offsets`, `num_items`, `num_segments` parameters.
For example, the following could be used to segment-sort a large buffer in two calls:
```
cub::DeviceSegmentedSort::SortKeys(
temp, temp_bytes, input, output, items/2, segments/2, offsets, offsets + 1);
```
This assumes items and segments are even value for illustration here.
The 2nd call would then just need a new offsets starting point
```
cub::DeviceSegmentedSort::SortKeys(
temp, temp_bytes, input, output, items/2, segments/2, offsets + (segments/2), offsets + (segments/2) + 1);
```
These offsetted offsets would read from the appropriate `input` positions and write to the appropriate `output` positions.

Unfortunately for significantly large values of `num_items` (~1M) the segmented-sort will read past the end of the temporary buffer.

Note the same error occurs with `cudf::DeviceSegmentedRadixSort::SortKeys`

### How to Reproduce

Here is a reproducer that shows the error
```
#include
#include
#include

#include
#include
#include
#include
#include

void segmented_sort(int const* input, int* output, int items, int segments, long const* offsets)
{
std::size_t temp_bytes = 0;
cub::DeviceSegmentedSort::SortKeys(
nullptr, temp_bytes, input, output, items, segments, offsets, offsets + 1);
void* temp;
cudaMalloc(&temp, temp_bytes);

printf("input(%p), output(%p), items=%d\n", input, output, items);
printf("offsets(%p), segments=%d\n", offsets, segments);
printf("temp(%p,%ld)\n", temp, temp_bytes);
std::cout << std::endl;

cub::DeviceSegmentedSort::SortKeys(
temp, temp_bytes, input, output, items, segments, offsets, offsets + 1);

auto rc = cudaStreamSynchronize(0);
std::cout << "segmented_sort = " << (int)rc << std::endl;
cudaFree(temp);
}

int main(int argc, const char** argv)
{
int N = 1'000'000'000;
if (argc > 1) { N = std::stoi(std::string(argv[1])); }
std::cout << "N = " << N << std::endl;

auto h_data = std::vector(N / 10);
auto d_data = thrust::device_vector(N);
for (int i = 0; i < 10; ++i) {
std::fill(h_data.begin(), h_data.end(), i);
cudaMemcpy(d_data.data().get() + (h_data.size() * i),
h_data.data(),
h_data.size() * sizeof(int),
cudaMemcpyHostToDevice);
}

// make 100 segments [0, 10M, 20M, 30M, ..., 1000M]
auto h_offsets = std::vector(100 + 1);
std::transform(thrust::counting_iterator(0),
thrust::counting_iterator(h_offsets.size()),
h_offsets.begin(),
[N](auto v) { return v * (N / 100L); });
auto d_offsets = thrust::device_vector(h_offsets.size());
cudaMemcpy(d_offsets.data().get(),
h_offsets.data(),
h_offsets.size() * sizeof(long),
cudaMemcpyHostToDevice);

auto d_output = thrust::device_vector(d_data.size());
auto input = d_data.data().get();
auto output = d_output.data().get();
auto offsets = d_offsets.data().get();

// segmented_sort(input, output, N, 100, offsets); // all of it (works)

// call segmented sort in 2 chunks
// segmented_sort(input, output, N / 2, 50, offsets); // 1st half (works)
segmented_sort(input, output, N / 2, 50, offsets + 50); // 2nd half (fails)

return 0;
}

```
This allocates too much memory to run in godbolt.
Use the following to compile (assuming source file name is segsort.cu):
```
nvcc -std=c++17 -o segsort segsort.cu
```
Run using no parameters
```
./segsort
```
Running with `compute-sanitizer` shows the error more specifically. Partial output:
```
compute-sanitizer --tool memcheck ./segsort
======== COMPUTE-SANITIZER
N = 1000000000
input(0x7fcfc8000000), output(0x7fced8000000), items=500000000
offsets(0x7fd0b6a00190), segments=50
temp(0x7fce60000000,2000000255)

========= Invalid __global__ write of size 4 bytes
========= at void cub::CUB_200400_520_NS::DeviceSegmentedSortFallbackKernel<(bool)0, cub::CUB_200400_520_NS::DeviceSegmentedSortPolicy::Policy860, int, cub::CUB_200400_520_NS::NullType, const long *, const long *, int>(const T3 *, T3 *, cub::CUB_200400_520_NS::detail::device_double_buffer, const T4 *, T4 *, cub::CUB_200400_520_NS::detail::device_double_buffer, T5, T6)+0x3b70
========= by thread (96,0,0) in block (0,0,0)
========= Address 0x7fced7359580 is out of bounds
========= and is 130 bytes after the nearest allocation at 0x7fce60000000 of size 2,000,000,255 bytes
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: [0x2c914f]
========= in /lib/x86_64-linux-gnu/libcuda.so.1
========= Host Frame:libcudart_static_4d8b33a106dceb3c07a56e26de61f2d53bb62a68 [0x58d2d]
========= in /home/dwendt/exper/cub/segsort/./segsort
========= Host Frame:cudaLaunchKernel [0xbc79d]
========= in /home/dwendt/exper/cub/segsort/./segsort
========= Host Frame:cudaError cudaLaunchKernel(char const*, dim3, dim3, void**, unsigned long, CUstream_st*) [0xf1d8]
========= in /home/dwendt/exper/cub/segsort/./segsort
========= Host Frame:__device_stub__ZN3cub17CUB_200400_520_NS33DeviceSegmentedSortFallbackKernelILb0ENS0_25DeviceSegmentedSortPolicyIiNS0_8NullTypeEE9Policy860EiS3_PKlS7_iEEvPKT1_PS8_NS0_6detail20device_double_bufferIS8_EEPKT2_PSF_NSD_ISF_EET3_T4_(int const*, int*, cub::CUB_200400_520_NS::detail::device_double_buffer&, cub::CUB_200400_520_NS::NullType const*, cub::CUB_200400_520_NS::NullType*, cub::CUB_200400_520_NS::detail::device_double_buffer&, long const*, long const*) [0xe744]
========= in /home/dwendt/exper/cub/segsort/./segsort
========= Host Frame:void cub::CUB_200400_520_NS::__wrapper__device_stub_DeviceSegmentedSortFallbackKernel::Policy860, int, cub::CUB_200400_520_NS::NullType, long const*, long const*, int>(int const*&, int*&, cub::CUB_200400_520_NS::detail::device_double_buffer&, cub::CUB_200400_520_NS::NullType const*&, cub::CUB_200400_520_NS::NullType*&, cub::CUB_200400_520_NS::detail::device_double_buffer&, long const*&, long const*&) [0xe7d0]
========= in /home/dwendt/exper/cub/segsort/./segsort
========= Host Frame:void cub::CUB_200400_520_NS::DeviceSegmentedSortFallbackKernel::Policy860, int, cub::CUB_200400_520_NS::NullType, long const*, long const*, int>(int const*, int*, cub::CUB_200400_520_NS::detail::device_double_buffer, cub::CUB_200400_520_NS::NullType const*, cub::CUB_200400_520_NS::NullType*, cub::CUB_200400_520_NS::detail::device_double_buffer, long const*, long const*) [0x4ad42]
....

```

### Expected behavior

Segmented sort should be able to sort in segments.

### Reproduction link

_No response_

### Operating System

Ubuntu Linux 20.04 and 22.04

### NVCC version
```
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Thu_Jun__6_02:18:23_PDT_2024
Cuda compilation tools, release 12.5, V12.5.82
Build cuda_12.5.r12.5/compiler.34385749_0
```

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

[BUG]: DeviceSegmentedSort fails with cudaErrorIllegalAddress when d_begin_offsets[0] != 0

Description

Contributor guide

Assessment

Get new issues in your inbox