[BUG]: thrust::min_element / max_element fail with numItems close to int_max

[fkallen]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Runtime Error

Component

Thrust

Describe the bug

terminate called after throwing an instance of 'thrust::system::system_error'
what(): extrema failed to synchronize: cudaErrorIllegalAddress: an illegal memory access was encountered
Aborted (core dumped)

I assume it happens because the sum in this loop overflows for int.

cccl/thrust/thrust/system/cuda/detail/reduce.h

Line 541 in 4d01371

while (block_offset + ITEMS_PER_TILE <= num_items)

How to Reproduce

#include <thrust/device_vector.h>
#include <thrust/extrema.h>

int main() {
  size_t N = std::numeric_limits<int>::max();
  thrust::device_vector<int> dArr(N, 1);
  thrust::max_element(dArr.begin(),dArr.end());
  return 0;
}

Expected behavior

No error

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[github-actions]

Hi @fkallen!

Thanks for submitting this issue - the CCCL team has been notified and we'll get back to you as soon as we can!
In the mean time, feel free to add any relevant information to this issue.

[gevtushenko]

Hello @fkallen and thank you for reporting this! Similar issue was addressed in CUB (1, 2). The thrust implementation is due to be replaced with the CUB one. This should address the issue. One note. The transition to CUB implementation is currently blocked by #548.