     CUDA_SAFE_CALL(cudaMemcpy(temp_buffer_->Get(), subgrid, 
                               subgrid_size.accumulate(rank) * elm_size,
                               cudaMemcpyHostToDevice));
-    
+
+    PSAssert(elm_size == 4 || elm_size == 8);
     if (rank == 3) {
-      Unpack3d<<<nblocks, thread_block_size>>>(
-          (int*)(temp_buffer_->Get()), (int*)Get(),
-          grid_size[0], grid_size[1], grid_size[2],
-          subgrid_offset[0], subgrid_offset[1], subgrid_offset[2],
-          subgrid_size[0], subgrid_size[1], subgrid_size[2]);
+      if (elm_size == 4) {
+        Unpack3d<float><<<nblocks, thread_block_size>>>(
+            (float*)(temp_buffer_->Get()), (float*)Get(),
+            grid_size[0], grid_size[1], grid_size[2],
+            subgrid_offset[0], subgrid_offset[1], subgrid_offset[2],
+            subgrid_size[0], subgrid_size[1], subgrid_size[2]);
+      } else if (elm_size == 8) {
+        Unpack3d<double><<<nblocks, thread_block_size>>>(
+            (double*)(temp_buffer_->Get()), (double*)Get(),
+            grid_size[0], grid_size[1], grid_size[2],
+            subgrid_offset[0], subgrid_offset[1], subgrid_offset[2],
+            subgrid_size[0], subgrid_size[1], subgrid_size[2]);
+      }        
     } else if (rank == 2) {
-      Unpack2d<<<nblocks, thread_block_size>>>(
-          (int*)(temp_buffer_->Get()), (int*)Get(), 
-          grid_size[0], grid_size[1], 
-          subgrid_offset[0], subgrid_offset[1],
-          subgrid_size[0], subgrid_size[1]);
+      if (elm_size == 4) {
+        Unpack2d<float><<<nblocks, thread_block_size>>>(
+            (float*)(temp_buffer_->Get()), (float*)Get(), 
+            grid_size[0], grid_size[1], 
+            subgrid_offset[0], subgrid_offset[1],
+            subgrid_size[0], subgrid_size[1]);
+      } else if (elm_size == 8) {
+        Unpack2d<double><<<nblocks, thread_block_size>>>(
+            (double*)(temp_buffer_->Get()), (double*)Get(), 
+            grid_size[0], grid_size[1], 
+            subgrid_offset[0], subgrid_offset[1],
+            subgrid_size[0], subgrid_size[1]);
+      }
     } else {
       LOG_ERROR() << "Unsupported\n";
       PSAbort(1);

                            cudaMemcpyDeviceToHost));
}

__global__ void Pack3D(const int *src, int *dst,
                       PSIndex xdim, PSIndex ydim, PSIndex zdim,
                       PSIndex xoff, PSIndex yoff, PSIndex zoff,
                       PSIndex xextent, PSIndex yextent, PSIndex zextent) {

  }
}

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