I'm having a confusing error when trying to compact multiple CUDA arrays using the thrust zip iterator. My case is simple: I have a thrust::vector of integers indicating the status of an object, and six vectors of floats indicating the object's position and velocity in 3D. When the the value of status (integer) vector element is 1, I want to remove the element from all vectors and shrink the vectors.
Following this example, I utilize a 7-way tuple and a Predicate.
// Predicate
struct isTupleFlagged
{
__host__ __device__ bool operator() ( const PartTuple& tup )
{
const int x = thrust::get<0>( tup );
return ( x == 1 );
}
};
void ParticleSystem::RemoveFlaggedParticles()
{
typedef thrust::device_vector< int >::iterator IntDIter;
typedef thrust::device_vector< float >::iterator FloatDIter;
typedef thrust::tuple< IntDIter, FloatDIter, FloatDIter, FloatDIter, FloatDIter, FloatDIter, FloatDIter > PartDIterTuple;
typedef thrust::zip_iterator< PartDIterTuple > ZipDIter;
ZipDIter newEnd = thrust::remove_if( thrust::make_zip_iterator( thrust::make_tuple( exit.begin(), x.begin(), y.begin(), z.begin(), vx.begin(), vy.begin(), vz.begin() ) ),
thrust::make_zip_iterator( thrust::make_tuple( exit.end(), x.end(), y.end(), z.end(), vx.end(), vy.end(), vz.end() ) ),
isTupleFlagged() );
PartDIterTuple endTuple = newEnd.get_iterator_tuple();
exit.erase(thrust::get<0>( endTuple ), exit.end());
x.erase( thrust::get<1>( endTuple ), x.end() );
y.erase( thrust::get<2>( endTuple ), y.end() );
z.erase( thrust::get<3>( endTuple ), z.end() );
vx.erase( thrust::get<4>( endTuple ), vx.end() );
vy.erase( thrust::get<5>( endTuple ), vy.end() );
vz.erase( thrust::get<6>( endTuple ), vz.end() );
}
However, this seems to erroneously remove extra elements. For example, when there is a single flagged object in a vector of 100, the post-compact size of the vectors is 52, not 99 as one would expect. What have I missed?