     thrust::for_each(input.begin(), input.end(), CopyFunctorWithManyRegisters<N>());
 }
 DECLARE_UNITTEST(TestForEachLargeRegisterFootprint);
+
+
+void TestForEachNLargeRegisterFootprint()
+{
+    const size_t N = 100;
+
+    thrust::device_vector<int> data(N, 12345);
+
+    thrust::device_vector<int *> input(1, thrust::raw_pointer_cast(&data[0])); // length is irrelevant
+    
+    thrust::for_each_n(input.begin(), input.size(), CopyFunctorWithManyRegisters<N>());
+}
+DECLARE_UNITTEST(TestForEachNLargeRegisterFootprint);
 
 
 template <typename T, unsigned int N>

DECLARE_VECTOR_UNITTEST(TestForEachSimple);


void TestForEachSimpleAnySpace(void)
{
    thrust::device_vector<int> output(7, 0);

    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;};
