     __host__ __device__
     bool operator()(T x) { return x ? true : false; }
 };
+
+template<typename T>
+struct mod_3
+{
+    __host__ __device__
+    unsigned int operator()(T x) { return static_cast<unsigned int>(x) % 3; }
+};
+    
     
 
 template <class Vector>

struct is_even
{
    __host__ __device__
    bool operator()(T x) { return (x & 1) == 0; }
};

template<typename T>

template <typename T>
void TestCopyIfStencil(const size_t n)
{
    thrust::host_vector<T>   h_data(n); thrust::sequence(h_data.begin(), h_data.end());
    thrust::device_vector<T> d_data(n); thrust::sequence(d_data.begin(), d_data.end()); 

    thrust::host_vector<T>   h_stencil = unittest::random_integers<T>(n);
    thrust::device_vector<T> d_stencil = unittest::random_integers<T>(n);

    thrust::host_vector<T>   h_result(n);
    thrust::device_vector<T> d_result(n);

    typename thrust::host_vector<T>::iterator   h_new_end;
    typename thrust::device_vector<T>::iterator d_new_end;

    // test with Predicate that returns a bool
    {
        thrust::host_vector<T>   h_result(n);
        thrust::device_vector<T> d_result(n);

        h_new_end = thrust::copy_if(h_data.begin(), h_data.end(), h_result.begin(), is_even<T>());
        d_new_end = thrust::copy_if(d_data.begin(), d_data.end(), d_result.begin(), is_even<T>());

        h_result.resize(h_new_end - h_result.begin());
        d_result.resize(d_new_end - d_result.begin());

        ASSERT_EQUAL(h_result, d_result);
    }
    
    // test with Predicate that returns a non-bool
    {
        thrust::host_vector<T>   h_result(n);
        thrust::device_vector<T> d_result(n);

        h_new_end = thrust::copy_if(h_data.begin(), h_data.end(), h_result.begin(), mod_3<T>());
        d_new_end = thrust::copy_if(d_data.begin(), d_data.end(), d_result.begin(), mod_3<T>());

        h_result.resize(h_new_end - h_result.begin());
        d_result.resize(d_new_end - d_result.begin());

        ASSERT_EQUAL(h_result, d_result);
    }
}
DECLARE_VARIABLE_UNITTEST(TestCopyIfStencil);


#if THRUST_DEVICE_BACKEND == THRUST_DEVICE_BACKEND_CUDA
// do we really want to test this ever?
