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

        {
        pdata_pos_tex.normalized = false;
        pdata_pos_tex.filterMode = cudaFilterModePoint;
        cudaBindTexture(0, pdata_pos_tex, d_pos, sizeof(Scalar4)*(N+nghosts));
        
        leaf_xyzf_tex.normalized = false;
        leaf_xyzf_tex.filterMode = cudaFilterModePoint;
        cudaBindTexture(0, leaf_xyzf_tex, d_leaf_xyzf, sizeof(Scalar4)*(N+nghosts));
        
        leaf_db_tex.normalized = false;
        leaf_db_tex.filterMode = cudaFilterModePoint;
        cudaBindTexture(0, leaf_db_tex, d_leaf_db, sizeof(Scalar2)*(N+nghosts));
        
        aabb_node_bounds_tex.normalized = false;
        aabb_node_bounds_tex.filterMode = cudaFilterModePoint;
        cudaBindTexture(0, aabb_node_bounds_tex, d_tree_aabbs, sizeof(Scalar4)*2*nnodes);
        
        head_list_tex.normalized = false;
        head_list_tex.filterMode = cudaFilterModePoint;
        cudaBindTexture(0, head_list_tex, d_head_list, sizeof(unsigned int)*(N+nghosts));
        }
    
    if (!filter_body && !diameter_shift)

    {
    cudaMemset(d_node_locks, 0, sizeof(unsigned int)*(nleafs - ntypes));
    
    gpu_nlist_bubble_aabbs_kernel<<<nleafs/block_size + 1, block_size>>>(d_node_locks,
                                                                         d_tree_aabbs,
                                                                         d_tree_parent_sib,
