         atomicMax(&d_conditions[type_i], n_neigh_i);
     }
     
-                                         
+/*!
+ * \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
+ * \param filter_body True if body filtering is enabled
+ * \param diameter_shift True if rcut(i,j) should be shifted by the particle diameters
+ * \param compute_capability Compute capability of the GPU (in 20, 30, 35 format)
+ * \param block_size Requested thread block size
+ *
+ * \returns cudaSuccess on completion
+ * \returns cudaError on failure to texture bind
+ *
+ * \note Kernel calls are templated on body filtering and diameter shifting for optimization.
+ * \note One thread is called for all leaf particles. Some of these threads will die because they correspond to ghost
+ *       particles not owned by the rank. Because the leaf particles are sorted, there is no easy way to skip these
+ *       particles, and this inefficiency is assumed to be relatively small.
+ */                                         
 cudaError_t gpu_nlist_traverse_tree(unsigned int *d_nlist,
                                     unsigned int *d_n_neigh,
                                     Scalar4 *d_last_updated_pos,

 * \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
 */
__global__ void gpu_nlist_move_particles_kernel(Scalar4 *d_leaf_xyzf,
                                                Scalar2 *d_leaf_db,


//! 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,
