Port the runtime_introspection functions to work in __host__ or __device__ code.
Port launch_calculator and helper code to work in __host__ or __device__ functions.
Annotate synchronize functions with __host__ __device__.
Port launch_closure to work in __host__ or __device__ code.
Fix some typos that were only visible in sm_35 code.
Make thrust::cuda::par visible in __device__ code.
Make thrust::for_each callable from __device__ code.
Make thrust::device visible in __device__ code.
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.
Eliminate superfluous #include <cassert>
Eliminate use of assert in throw_on_error. Generate a trap instruction instead.

We may want to halt kernel execution a different way in the future, though.
Move terminate into its own header.
Eliminate use of assert() in temporary_allocator.
Test gather with thrust::device in __device__ code.
Test replace algorithms with thrust::device in __device__ code.
Test scatter algorithms with thrust::device in __device__ code.