  * \param d_body Particle body ids
  * \param d_map_tree_pid ParticleData indexes corresponding to a leaf particle id
  * \param Ntot Number of particles owned by this rank
+ *
+ * \b Implementation
+ * One thread per particle is called. Writes are coalesced by writing in leaf order, and reading in a scattered way.
  */
 __global__ void gpu_nlist_move_particles_kernel(Scalar4 *d_leaf_xyzf,
                                                 Scalar2 *d_leaf_db,

    {
    cudaMemset(d_node_locks, 0, sizeof(unsigned int)*(nleafs - ntypes));
    
    // later we can use block size to perform atomic operations in shared memory rather than
    // global memory for nodes in the same block. but for now, just use slow global memory for everything.
    gpu_nlist_bubble_aabbs_kernel<<<nleafs/block_size + 1, block_size>>>(d_node_locks,
                                                                         d_tree_aabbs,
                                                                         d_tree_parent_sib,

 * \b Implementation
 * One thread is called per leaf node. The second thread to reach an internal node processes its two children,
 * which guarantees that no node AABB is prematurely processed. The arrival order at a node is controlled by an atomic
 * thread lock in global memory. This locking could be accelerated by using shared memory whenever a node is being
 * processed by threads in the same block.
 *
 * When processing the node, the thread also walks up the tree to find the "rope" that tells a traverser
 * how to navigate the tree. If a query AABB intersects the current node, then the traverser always moves the the left
