How to prevent thrust::copy_if illegal memory access

Question

In a CUDA C++ code, I am using thrust::copy_if to copy those integer values that are not -1 from the array x to y:

thrust::copy_if(x_ptr, x_ptr + N, y_ptr, is_not_minus_one());

I put the code in a try/catch and I update the array x regularly, and again push numbers that are not -1 to the end of the y until it accessed to the out of the array and returns

CUDA Runtime Error: an illegal memory access was encountered

As I don't know how many values are not -1 in each iteration, I should keep generating and copying numbers in the main array until it becomes full. However, I wanna it to stop once it reaches the end of the array. How can I manage it?

A naive idea would be using another array and then copying if the number of new values does not exceed the size. but it might be inefficient. Any better idea?

Answer 1

Here is one possible method:

use remove_if() instead, (with the opposite condition - remove if -1) on x. Then use (x start and) the returned iterator to copy to your final array (y). Every time you do a copy to y, you will know how much space is left in y, and you can adjust the copy quantity if needed, before doing the copy.

Example (removing 1 values):

$ cat t2088.cu
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/complex.h>
#include <thrust/transform.h>
#include <thrust/copy.h>
#include <thrust/functional.h>
#include <iostream>

using mt = int;
using namespace thrust::placeholders;
int main(){

  thrust::device_vector<mt> a(5);
  thrust::device_vector<mt> c(25);
  bool done = false;
  int empty_spaces = c.size();
  int copy_start_index = 0;
  while (!done){
    thrust::sequence(a.begin(), a.end());
    int num_to_copy = (thrust::remove_if(a.begin(), a.end(), _1 == 1) - a.begin());
    int actual_num_to_copy = std::min(num_to_copy, empty_spaces);
    if (actual_num_to_copy != num_to_copy) done = true;
    thrust::copy_n(a.begin(), actual_num_to_copy, c.begin()+copy_start_index);
    copy_start_index += actual_num_to_copy;
    empty_spaces -= actual_num_to_copy;
    }
  std::cout << "output array is full!" << std::endl;
  thrust::host_vector<mt> h_c = c;
  thrust::copy(h_c.begin(), h_c.end(), std::ostream_iterator<mt>(std::cout, ","));
  std::cout << std::endl;
}
$ nvcc -o t2088 t2088.cu
$ compute-sanitizer ./t2088
========= COMPUTE-SANITIZER
output array is full!
0,2,3,4,0,2,3,4,0,2,3,4,0,2,3,4,0,2,3,4,0,2,3,4,0,
========= ERROR SUMMARY: 0 errors
$