thrust
thrust copied to clipboard
copy_if integer overflow for sizes > 2^31
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.
We should triage this by using THRUST_INDEX_TYPE_DISPATCH .
Related to #1271.
I am seeing a similar issue with thrust::remove & thrust::remove_if. Sounds related and better be fixed together.