thrust partition fails for size_t{3} * 1024 * 1024 * 1024 elements

[seunghwak]

Test environment
GA102 (48 GB GPU memory)
CUDA 11.4
Thrust 1.15.0

The code below fails with cudaErrorIllegalAddress

It seems like this works up to size_t{2} * 1024 * 1024 * 1024 + small number (didn't crash, haven't verified the results), but fails with size_t{3} * 1024 * 1024 * 1024 and size_t{4} * 1024 * 1024 * 1024. I suspect some sort of an integer overflow issue.

  thrust::device_vector<int8_t> values(size_t{3} * 1024 * 1024 * 1024);
  thrust::tabulate(thrust::device, values.begin(), values.end(), []__device__(auto i) { return static_cast<int8_t>(i % 2); });
  thrust::partition(thrust::device, values.begin(), values.end(), []__device__(auto v) { return (v == int8_t{0}); });

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  CUDA free failed: cudaErrorIllegalAddress: an illegal memory access was encountered

[lilohuang]

I think it's a known limitation. NVIDIA/cccl#744

[alliepiper]

@lilohuang is correct -- this is a known issue that's tracked by NVIDIA/cub#212.

[ericniebler]

Closed as duplicate