 }
 DECLARE_UNITTEST(TestForEachWithLargeTypes);
 
+
+template <typename T, unsigned int N>
+void _TestForEachNWithLargeTypes(void)
+{
+    size_t n = (64 * 1024) / sizeof(FixedVector<T,N>);
+
+    thrust::host_vector< FixedVector<T,N> > h_data(n);
+
+    for(size_t i = 0; i < h_data.size(); i++)
+        h_data[i] = FixedVector<T,N>(i);
+
+    thrust::device_vector< FixedVector<T,N> > d_data = h_data;
+   
+    SetFixedVectorToConstant<T,N> func(123);
+
+    thrust::for_each_n(h_data.begin(), h_data.size(), func);
+    thrust::for_each_n(d_data.begin(), d_data.size(), func);
+
+    ASSERT_EQUAL_QUIET(h_data, d_data);
+}
+
+
+void TestForEachNWithLargeTypes(void)
+{
+    _TestForEachNWithLargeTypes<int,    1>();
+    _TestForEachNWithLargeTypes<int,    2>();
+    _TestForEachNWithLargeTypes<int,    4>();
+    _TestForEachNWithLargeTypes<int,    8>();
+    _TestForEachNWithLargeTypes<int,   16>();
+
+    KNOWN_FAILURE;
+
+    //_TestForEachNWithLargeTypes<int,   32>();  // fails on Linux 32 w/ gcc 4.1
+    //_TestForEachNWithLargeTypes<int,   64>();
+    //_TestForEachNWithLargeTypes<int,  128>();
+    //_TestForEachNWithLargeTypes<int,  256>();
+    //_TestForEachNWithLargeTypes<int,  512>();
+    //_TestForEachNWithLargeTypes<int, 1024>();  // fails on Vista 64 w/ VS2008
+}
+DECLARE_UNITTEST(TestForEachNWithLargeTypes);
+
 __THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING_END

    thrust::for_each(input.begin(), input.end(), CopyFunctorWithManyRegisters<N>());
}
DECLARE_UNITTEST(TestForEachLargeRegisterFootprint);


template <typename T, unsigned int N>

    ASSERT_EQUAL_QUIET(d_result, d_input.end());
}
DECLARE_VARIABLE_UNITTEST(TestForEach);


template <typename T>
void TestForEachN(const size_t n)
{
    const size_t output_size = std::min((size_t) 10, 2 * n);
    
    thrust::host_vector<T> h_input = unittest::random_integers<T>(n);

    for(size_t i = 0; i < n; i++)
        h_input[i] =  ((size_t) h_input[i]) % output_size;
    
    thrust::device_vector<T> d_input = h_input;

    thrust::host_vector<T>   h_output(output_size, (T) 0);
    thrust::device_vector<T> d_output(output_size, (T) 0);

    mark_present_for_each<T> h_f;
    mark_present_for_each<T> d_f;
    h_f.ptr = &h_output[0];
    d_f.ptr = (&d_output[0]).get();
    
    typename thrust::host_vector<T>::iterator h_result =
      thrust::for_each_n(h_input.begin(), h_input.size(), h_f);

    typename thrust::device_vector<T>::iterator d_result =
      thrust::for_each_n(d_input.begin(), d_input.size(), d_f);

    ASSERT_EQUAL(h_output, d_output);
    ASSERT_EQUAL_QUIET(h_result, h_input.end());
    ASSERT_EQUAL_QUIET(d_result, d_input.end());
}
DECLARE_VARIABLE_UNITTEST(TestForEachN);


template <size_t N> __host__ __device__ void f   (int * x) { int temp = *x; f<N - 1>(x + 1); *x = temp;};
