   uint bid = blockIdx.y * gridDim.x + blockIdx.x;
   uint tid = threadIdx.x;
   uint id  = bid * blockDim.x + tid;
+
+
 #if 0
   //slowest
   if(id >= n_extract) return;

    } //end z
  }  //end y
} //end x, function



typedef struct bodyStruct
{
  real4  acc0;
  real4  Ppos;
  real4  Pvel;
  float2 time;
  unsigned long long id;

#ifdef DO_BLOCK_TIMESTEP_EXCHANGE_MPI
  uint4 key;
  real4 pos;
  real4 vel;
  real4 acc1;
#endif
} bodyStruct;




#endif






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