Don't use the __CUDA_ARCH__ workaround in launch_closure.
Use cudaDeviceSynchronize instead of cudaThreadSynchronize.
Make annotate uninitialized and uninitialized_array's member functions with __host__ __device__.
Use uninitialized to marshal kernel parameters which are passed by value for compatibility with kernels launched from __device__ code.
Eliminate a warning in use_big_closure().
Use the __host__ __device__ workaround technique in for_each_n.
Test for_each & for_each_n with thrust::device in __device__ code.