     return cudaSuccess;
     }
 
+//! Kernel to find divisons between particle types in sorted order
+/*!
+ * \param d_type_head Index to first type in leaf ordered particles by type
+ * \param d_pos Particle positions
+ * \param d_map_tree_pid ParticleData indexes corresponding to a leaf particle id
+ * \param N Total number of particles on rank (including ghosts)
+ *
+ * The starting index for each type of particles is the first particle where the left neighbor is not of the same type.
+ * One thread per particle is needed to check this, skipping the thread for the first particle because 0 is trivially
+ * the head index of the first type.
+ */
 __global__ void gpu_nlist_get_divisions_kernel(unsigned int *d_type_head,
                                                const Scalar4 *d_pos,
                                                const unsigned int *d_map_tree_pid,

 * \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.
 *
 * 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


//! Kernel for traversing tree to generate neighbor list
/*!
 * \param d_nlist Neighbor list for writing
 * \param d_n_neigh Number of neighbors per particle
 * \param d_last_updated_pos Records current particle positions
 * \param d_conditions Store overflow condition by type
 * \param d_Nmax Maximum number of neighbors allocated by type
 * \param d_head_list Indexes for writing into neighbor list
 * \param N Number of particles
 * \param nghosts Number of ghost particles
 * \param d_map_tree_pid Map leaf index to local particle index
 * \param d_leaf_offset Offset for reading leaf particles by type
 * \param d_tree_roots Index for tree root by type
 * \param d_tree_aabbs Tree AABBs
 * \param nleafs Total number of leafs
 * \param d_leaf_xyzf Leaf position-id array
 * \param d_leaf_db Leaf diameter-body array
 * \param d_pos Particle positions
 * \param d_image_list Translation vectors to check for traversal
 * \param nimages Number of translation vectors to check
 * \param d_r_cut Cutoff radius by type r_cut(i,j)
 * \param r_buff Buffer around cutoff radius
 * \param max_diam Maximum diameter attained by a particle for diameter shifting
 * \param ntypes Number of particle types
 *
 * \b Implementation
 * One thread is launched per particle, but the threads operate on particles in leaf order rather than ParticleData
 * order in order to minimize divergence within a warp (particles in the same leaf should intersect similar parts of the
 * tree). Each thread iterates on the particle types (trees) and queries on all translation vectors using a stackless
 * search. When the query AABB intersects a node AABB, the node AABB is checked to be an internal node or a leaf node.
 * If an internal node, then the traversal advances to that node's left child. If a leaf node, the leaf particles are
 * tested directly to be included in the neighbor list. The node then advances along that leaf node's rope. If the AABB
 * is not intersected, the traversal advances along the rope. This process proceeds until a rope signals that the
 * traversal is complete.
 */
template<unsigned char flags>
__global__ void gpu_nlist_traverse_tree_kernel(unsigned int *d_nlist,
