thrust icon indicating copy to clipboard operation
thrust copied to clipboard

copy_if integer overflow for sizes > 2^31

Open RAMitchell opened this issue 5 years ago • 3 comments

For sizes > 2^31 we get an integer overflow with copy_if using the device backend. This results in out of memory:

terminate called after throwing an instance of 'thrust::system::detail::bad_alloc'
  what():  std::bad_alloc: temporary_buffer::allocate: get_temporary_buffer failed

MRE:

#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <limits>

int main() {
  // size_t n = std::numeric_limits<int>::max() - 1000ll; // Works
  size_t n = std::numeric_limits<int>::max() + 1ll; // Error
  std::cout << "n size: " << n << std::endl;
  thrust::device_vector<int> v(n);
  thrust::device_vector<int> result(n);
  thrust::copy_if(v.begin(), v.end(), result.begin(), thrust::identity<int>());
}

Suspected cause: https://github.com/NVIDIA/thrust/blob/e4d96a2ecaae1fb2964be8caace289e3c314ac7b/thrust/system/cuda/detail/copy_if.h#L704 Custom cub-like backend, still using int size type.

RAMitchell avatar Oct 05 '20 04:10 RAMitchell

We should triage this by using THRUST_INDEX_TYPE_DISPATCH .

alliepiper avatar Oct 05 '20 16:10 alliepiper

Related to #1271.

alliepiper avatar Oct 13 '20 16:10 alliepiper

I am seeing a similar issue with thrust::remove & thrust::remove_if. Sounds related and better be fixed together.

seunghwak avatar May 10 '22 18:05 seunghwak