                             cudaMemcpyDeviceToHost));
 }
 
-__global__ void Pack3D(const int *src, int *dst,
+//! Pack a 3-dimensional sub grid into a linear region
+/*!
+  \param src 3D source grid 
+  \param dst Linear buffer
+  \param xdim X size of the grid
+  \param ydim Y size of the grid
+  \param zdim Z size of the grid
+  \param xoff X offset of the sub grid
+  \param yoff Y offset of the sub grid
+  \param zoff Z offset of the sub grid
+  \param xextent X size of the sub grid
+  \param yextent Y size of the sub grid
+  \param zextent Z size of the sub grid 
+ */
+template <class T>
+__global__ void Pack3D(const T *src, T *dst,
                        PSIndex xdim, PSIndex ydim, PSIndex zdim,
                        PSIndex xoff, PSIndex yoff, PSIndex zoff,
                        PSIndex xextent, PSIndex yextent, PSIndex zextent) {

                            cudaMemcpyHostToDevice));
}

__global__ void Unpack2d(const int *src, int *dst,
                         PSIndex xdim, PSIndex ydim,
                         PSIndex xoff, PSIndex yoff,
                         PSIndex xextent, PSIndex yextent) {

  CUDA_SAFE_CALL(cudaFree(ptr));
}

template <class T>
__global__ void Pack2D(const T *src, T *dst,
                       PSIndex xdim, PSIndex ydim,
                       PSIndex xoff, PSIndex yoff,
                       PSIndex xextent, PSIndex yextent) {
