[BUG]: thrust::merge_by_key crashes on large data blocks
- 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
Thrust
### Describe the bug
If I incrementally add high-number arrays to thrust::merge_by_key it crashes even after having enough memory remaining. But on the same-sized arrays, thrust::sort doesn't fail.
```
// timestamps has int64_t type data.
HOST void temporal_graph::sort_and_merge_edges_cuda(TemporalGraphStore *graph, const size_t start_idx) {
const size_t total_size = edge_data::size(graph->edge_data);
const size_t new_edges_count = total_size - start_idx;
if (new_edges_count == 0) return;
int* d_sources = graph->edge_data->sources;
int* d_targets = graph->edge_data->targets;
int64_t* d_timestamps = graph->edge_data->timestamps;
// === Step 1: Create index array for new edges ===
size_t* d_indices = nullptr;
CUDA_CHECK_AND_CLEAR(cudaMalloc(&d_indices, new_edges_count * sizeof(size_t)));
thrust::sequence(DEVICE_EXECUTION_POLICY, d_indices, d_indices + new_edges_count, start_idx);
// === Step 2: Sort new edge indices by timestamp ===
thrust::sort(
DEVICE_EXECUTION_POLICY,
d_indices, d_indices + new_edges_count,
[d_timestamps] __device__(size_t i, size_t j) {
return d_timestamps[i] < d_timestamps[j];
});
CUDA_KERNEL_CHECK("After thrust sort");
// === Step 3: Allocate temporary arrays for sorted edges ===
int *d_sorted_sources = nullptr, *d_sorted_targets = nullptr;
int64_t *d_sorted_timestamps = nullptr;
CUDA_CHECK_AND_CLEAR(cudaMalloc(&d_sorted_sources, new_edges_count * sizeof(int)));
CUDA_CHECK_AND_CLEAR(cudaMalloc(&d_sorted_targets, new_edges_count * sizeof(int)));
CUDA_CHECK_AND_CLEAR(cudaMalloc(&d_sorted_timestamps, new_edges_count * sizeof(int64_t)));
// === Step 4: Gather new edges in sorted order ===
thrust::gather(DEVICE_EXECUTION_POLICY, d_indices, d_indices + new_edges_count, d_sources, d_sorted_sources);
thrust::gather(DEVICE_EXECUTION_POLICY, d_indices, d_indices + new_edges_count, d_targets, d_sorted_targets);
thrust::gather(DEVICE_EXECUTION_POLICY, d_indices, d_indices + new_edges_count, d_timestamps, d_sorted_timestamps);
CUDA_KERNEL_CHECK("After gather");
clear_memory(&d_indices, true);
// === Step 5: Allocate merged output arrays ===
int *d_merged_sources = nullptr, *d_merged_targets = nullptr;
int64_t* d_merged_timestamps = nullptr;
CUDA_CHECK_AND_CLEAR(cudaMalloc(&d_merged_sources, total_size * sizeof(int)));
CUDA_CHECK_AND_CLEAR(cudaMalloc(&d_merged_targets, total_size * sizeof(int)));
CUDA_CHECK_AND_CLEAR(cudaMalloc(&d_merged_timestamps, total_size * sizeof(int64_t)));
// === Step 6: Merge using zipped source/target tuple ===
using ZipIterator = thrust::zip_iterator>;
using ConstZipIterator = thrust::zip_iterator>;
ConstZipIterator old_edges = thrust::make_zip_iterator(thrust::make_tuple(d_sources, d_targets));
ConstZipIterator new_edges = thrust::make_zip_iterator(thrust::make_tuple(d_sorted_sources, d_sorted_targets));
ZipIterator merged_edges = thrust::make_zip_iterator(thrust::make_tuple(d_merged_sources, d_merged_targets));
// CRASHES HERE
thrust::merge_by_key(
DEVICE_EXECUTION_POLICY,
d_timestamps, d_timestamps + start_idx,
d_sorted_timestamps, d_sorted_timestamps + new_edges_count,
old_edges, new_edges,
d_merged_timestamps, merged_edges,
thrust::less());
CUDA_KERNEL_CHECK("After zipped merge_by_key");
// === Step 7: Copy back merged results ===
CUDA_CHECK_AND_CLEAR(cudaMemcpy(d_sources, d_merged_sources, total_size * sizeof(int), cudaMemcpyDeviceToDevice));
CUDA_CHECK_AND_CLEAR(cudaMemcpy(d_targets, d_merged_targets, total_size * sizeof(int), cudaMemcpyDeviceToDevice));
CUDA_CHECK_AND_CLEAR(cudaMemcpy(d_timestamps, d_merged_timestamps, total_size * sizeof(int64_t), cudaMemcpyDeviceToDevice));
// === Step 8: Cleanup ===
clear_memory(&d_sorted_sources, true);
clear_memory(&d_sorted_targets, true);
clear_memory(&d_sorted_timestamps, true);
clear_memory(&d_merged_sources, true);
clear_memory(&d_merged_targets, true);
clear_memory(&d_merged_timestamps, true);
}
```
This is part of a bigger application that's exposed to Python using Pybind.
Now when I incrementally add a higher number of data instances, this crashes.
```
(venv) (base) ../cuda-12.6/compute-sanitizer/compute-sanitizer python test_sliding_window.py --use_gpu
========= COMPUTE-SANITIZER
Loaded 63,497,050 edges.
---- Running on GPU. ----
--- Adding edges 0 to 5,000,000 ---
Edge addition time: 3.465 sec
Walk sampling time: 1.318 sec
Active edges in graph: 4,092,146 (with sliding window)
--- Adding edges 5,000,000 to 10,000,000 ---
Edge addition time: 4.520 sec
Walk sampling time: 1.276 sec
Active edges in graph: 6,341,331 (with sliding window)
--- Adding edges 10,000,000 to 15,000,000 ---
Edge addition time: 5.035 sec
Walk sampling time: 1.258 sec
Active edges in graph: 7,178,464 (with sliding window)
--- Adding edges 15,000,000 to 20,000,000 ---
Edge addition time: 4.986 sec
Walk sampling time: 1.254 sec
Active edges in graph: 7,077,491 (with sliding window)
--- Adding edges 20,000,000 to 25,000,000 ---
Edge addition time: 5.023 sec
Walk sampling time: 1.250 sec
Active edges in graph: 7,077,491 (with sliding window)
--- Adding edges 25,000,000 to 30,000,000 ---
Edge addition time: 4.973 sec
Walk sampling time: 1.256 sec
Active edges in graph: 7,077,491 (with sliding window)
--- Adding edges 30,000,000 to 35,000,000 ---
Edge addition time: 4.973 sec
Walk sampling time: 1.253 sec
Active edges in graph: 7,077,491 (with sliding window)
--- Adding edges 35,000,000 to 40,000,000 ---
Edge addition time: 4.973 sec
Walk sampling time: 1.256 sec
Active edges in graph: 7,077,491 (with sliding window)
--- Adding edges 40,000,000 to 45,000,000 ---
Edge addition time: 6.921 sec
Walk sampling time: 1.272 sec
Active edges in graph: 10,038,645 (with sliding window)
--- Adding edges 45,000,000 to 50,000,000 ---
Edge addition time: 10.018 sec
Walk sampling time: 1.278 sec
Active edges in graph: 15,036,976 (with sliding window)
--- Adding edges 50,000,000 to 55,000,000 ---
========= Invalid __shared__ read of size 8 bytes
========= at void thrust::THRUST_200500_750_NS::cuda_cub::core::_kernel_agent>, thrust::THRUST_200500_750_NS::zip_iterator>, long, long *, thrust::THRUST_200500_750_NS::zip_iterator>, thrust::THRUST_200500_750_NS::less, cuda::std::__4::integral_constant>, long *, long *, thrust::THRUST_200500_750_NS::zip_iterator>, thrust::THRUST_200500_750_NS::zip_iterator>, long, long, long *, thrust::THRUST_200500_750_NS::zip_iterator>, thrust::THRUST_200500_750_NS::less, long *>(T2, T3, T4, T5, T6, T7, T8, T9, T10, T11)+0x1440
========= by thread (0,0,0) in block (6998,0,0)
========= Address 0xffffe7c0 is out of bounds
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: [0x2dfbef]
========= in /lib64/libcuda.so.1
========= Host Frame: [0x15aa7]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/temporal_random_walk.libs/libcudart-09529672.so.12.6.77
========= Host Frame:cudaLaunchKernel [0x759f0]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/temporal_random_walk.libs/libcudart-09529672.so.12.6.77
========= Host Frame: [0x9630a]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x9a124]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x9b966]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x1da28]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x437b5]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame:cfunction_call in /usr/local/src/conda/python-3.9.13/Objects/methodobject.c:543 [0x108126]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyObject_MakeTpCall in /usr/local/src/conda/python-3.9.13/Objects/call.c:191 [0xf0edb]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:method_vectorcall in /usr/local/src/conda/python-3.9.13/Objects/classobject.c:53 [0x10604f]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.9.13/Python/ceval.c:3489 [0xeccee]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:function_code_fastcall in /usr/local/src/conda/python-3.9.13/Objects/call.c:330 [0xf87f2]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.9.13/Python/ceval.c:3520 [0xe83a0]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:function_code_fastcall in /usr/local/src/conda/python-3.9.13/Objects/call.c:330 [0xf87f2]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.9.13/Python/ceval.c:3520 [0xe83a0]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalCode in /usr/local/src/conda/python-3.9.13/Python/ceval.c:4329 [0xe70c9]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalCodeWithName in /usr/local/src/conda/python-3.9.13/Python/ceval.c:4361 [0xe6d56]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:PyEval_EvalCodeEx in /usr/local/src/conda/python-3.9.13/Python/ceval.c:4377 [0xe6d08]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:PyEval_EvalCode in /usr/local/src/conda/python-3.9.13/Python/ceval.c:828 [0x194e7a]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:run_eval_code_obj in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:1221 [0x1c2306]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:run_mod in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:1242 [0x1be26f]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:pyrun_file.cold in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:1140 [0x563ec]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:PyRun_SimpleFileExFlags in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:483 [0x1b8061]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:Py_RunMain in /usr/local/src/conda/python-3.9.13/Modules/main.c:683 [0x1b55cd]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:Py_BytesMain in /usr/local/src/conda/python-3.9.13/Modules/main.c:1129 [0x188ff8]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:__libc_start_call_main [0x295cf]
========= in /lib64/libc.so.6
========= Host Frame:__libc_start_main [0x2967f]
========= in /lib64/libc.so.6
========= Host Frame: [0x188ead]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
=========
========= Invalid __shared__ read of size 8 bytes
========= at void thrust::THRUST_200500_750_NS::cuda_cub::core::_kernel_agent>, thrust::THRUST_200500_750_NS::zip_iterator>, long, long *, thrust::THRUST_200500_750_NS::zip_iterator>, thrust::THRUST_200500_750_NS::less, cuda::std::__4::integral_constant>, long *, long *, thrust::THRUST_200500_750_NS::zip_iterator>, thrust::THRUST_200500_750_NS::zip_iterator>, long, long, long *, thrust::THRUST_200500_750_NS::zip_iterator>, thrust::THRUST_200500_750_NS::less, long *>(T2, T3, T4, T5, T6, T7, T8, T9, T10, T11)+0x1440
========= by thread (1,0,0) in block (6998,0,0)
========= Address 0xffffe7d8 is out of bounds
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: [0x2dfbef]
========= in /lib64/libcuda.so.1
========= Host Frame: [0x15aa7]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/temporal_random_walk.libs/libcudart-09529672.so.12.6.77
========= Host Frame:cudaLaunchKernel [0x759f0]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/temporal_random_walk.libs/libcudart-09529672.so.12.6.77
========= Host Frame: [0x9630a]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x9a124]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x9b966]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x1da28]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame: [0x437b5]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/lib/python3.9/site-packages/_temporal_random_walk.cpython-39-x86_64-linux-gnu.so
========= Host Frame:cfunction_call in /usr/local/src/conda/python-3.9.13/Objects/methodobject.c:543 [0x108126]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyObject_MakeTpCall in /usr/local/src/conda/python-3.9.13/Objects/call.c:191 [0xf0edb]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:method_vectorcall in /usr/local/src/conda/python-3.9.13/Objects/classobject.c:53 [0x10604f]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.9.13/Python/ceval.c:3489 [0xeccee]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:function_code_fastcall in /usr/local/src/conda/python-3.9.13/Objects/call.c:330 [0xf87f2]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.9.13/Python/ceval.c:3520 [0xe83a0]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:function_code_fastcall in /usr/local/src/conda/python-3.9.13/Objects/call.c:330 [0xf87f2]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.9.13/Python/ceval.c:3520 [0xe83a0]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalCode in /usr/local/src/conda/python-3.9.13/Python/ceval.c:4329 [0xe70c9]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:_PyEval_EvalCodeWithName in /usr/local/src/conda/python-3.9.13/Python/ceval.c:4361 [0xe6d56]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:PyEval_EvalCodeEx in /usr/local/src/conda/python-3.9.13/Python/ceval.c:4377 [0xe6d08]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:PyEval_EvalCode in /usr/local/src/conda/python-3.9.13/Python/ceval.c:828 [0x194e7a]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:run_eval_code_obj in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:1221 [0x1c2306]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:run_mod in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:1242 [0x1be26f]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:pyrun_file.cold in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:1140 [0x563ec]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:PyRun_SimpleFileExFlags in /usr/local/src/conda/python-3.9.13/Python/pythonrun.c:483 [0x1b8061]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:Py_RunMain in /usr/local/src/conda/python-3.9.13/Modules/main.c:683 [0x1b55cd]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:Py_BytesMain in /usr/local/src/conda/python-3.9.13/Modules/main.c:1129 [0x188ff8]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
========= Host Frame:__libc_start_call_main [0x295cf]
========= in /lib64/libc.so.6
========= Host Frame:__libc_start_main [0x2967f]
========= in /lib64/libc.so.6
========= Host Frame: [0x188ead]
========= in /mnt/nfs2/inf/ms2420/temporal-random-walk-test/venv/bin/python
```
If I modify this to do a full-sort (which I am having to do temporarily), then I get no issue, but doing full-sort on two large arrays isn't very efficient.
My question is, are you aware of this? I am getting this on Cuda 12.6, if I use more recent Cuda driver, will I get the same issue there too? I didn't try that way yet because I am on an HPC environment with no sudo access.
Or am I doing anything wrong?
### How to Reproduce
Calling thrust::merge_by_key with more than 30 million sized int64_t data blocks.
### Expected behavior
Successful merge, regardless of size, as long as there is free memory available.
### Reproduction link
_No response_
### Operating System
Rockylinux 9
### nvidia-smi output
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 560.35.05 Driver Version: 560.35.05 CUDA Version: 12.6 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA A40 On | 00000000:21:00.0 Off | 0 |
| 0% 32C P8 21W / 300W | 1MiB / 46068MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| No running processes found |
+-----------------------------------------------------------------------------------------+
### NVCC version
Cuda 12.6
Contributor guide
Assessment
This issue has not been assessed yet.