 
   if(id < n_extract)
   {
+	shmem[threadIdx.x].pos   = pos[extractList[offset+id].y];
+	shmem[threadIdx.x].vel   = vel[extractList[offset+id].y];
     shmem[threadIdx.x].Ppos  = Ppos[extractList[offset+id].y];
     shmem[threadIdx.x].Pvel  = Pvel[extractList[offset+id].y];
     shmem[threadIdx.x].acc0  = acc0[extractList[offset+id].y];
     shmem[threadIdx.x].time  = time[extractList[offset+id].y];
 
-    shmem[threadIdx.x].id    = body_id[extractList[offset+id].y];
+    shmem[threadIdx.x].id     = body_id[extractList[offset+id].y];
     shmem[threadIdx.x].Pvel.w = h[extractList[offset+id].y];
 
 
 #ifdef DO_BLOCK_TIMESTEP_EXCHANGE_MPI
     shmem[threadIdx.x].key   = body_key[extractList[offset+id].y];
-    shmem[threadIdx.x].pos   = pos[extractList[offset+id].y];
-    shmem[threadIdx.x].vel   = vel[extractList[offset+id].y];
     shmem[threadIdx.x].acc1  = acc1[extractList[offset+id].y];
 #endif
   }

  destination[id].id    = body_id[extractList[offset+id].y];
  destination[id].key   = body_key[extractList[offset+id].y];

#elif 1
  //5x faster than original (above)
  __shared__ bodyStruct shmem[128];


#endif






KERNEL_DECLARE(gpu_internalMoveSFC) (int       n_extract,
                                  int       n_bodies,
                                  uint4  lowBoundary,
