How to handle an output of uncertain size in CUDA

I have a large array that I want to put in CUDA for some conditional evaluations, and I want to output the values that satisfy the conditions. However, before the evaluations, I don't know the total number of output values. I only know that it will be less than 30% of the total count. Therefore, I want to know how to write the code to minimize memory usage as much as possible.

To put my comment into code, this is how it could be done using Thrust.

#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/count.h>
#include <iostream>
#include <cassert>
struct IsEven{
    __host__ __device__ 
    bool operator()(int i) const{
        return i % 2 == 0;
    }
};
int main(){
    //fill data with 0,1,2,3,4
    thrust::device_vector<int> d_data(5);
    thrust::sequence(d_data.begin(), d_data.end(), 0);
    //count number of even elements
    int numEvenElements = thrust::count_if(d_data.begin(), d_data.end(), IsEven{});
    //allocate memory
    thrust::device_vector<int> d_evenElements(numEvenElements);
    //copy even elements
    auto copyend = thrust::copy_if(d_data.begin(), d_data.end(), d_evenElements.begin(), IsEven{});
    assert(copyend == d_evenElements.end());
    
    for(int i = 0; i < numEvenElements; i++){
        std::cout << d_evenElements[i] << "\n";
    }
}

In general, to count the number of elements satisfying a condition, you would want to create a range where for each input element you put 1 if the condition is true and 0 otherwise. Then, simply perform a parallel sum reduction over this range. With fancy iterators like a thrust::transform_iterator, the range does not need to be materialized in memory. thrust::count_if does all that for you.

There are other approaches that depend on how much memory you actually have available:

One other approach to avoid an output buffer that is not mostly empty would be to allocate 100% of the input size. Compact the selected elements into the large buffer, then copy them to the small buffer of exact size and free the large buffer. Of course, this fails if one of the buffers cannot be allocated.

A third approach would be to use the virtual memory driver API to allocate a 30% buffer, and manually release the unoccupied memory pages after compaction.