 
         // now that possible neighbor checks are finished, done (for the cta) depends only on first thread
         // neighbor list only needs to get written into if thread 0 is not done
-        all_done = hoomd::detail::WarpScan<bool, threads_per_particle>().Broadcast(this_thread_done, 0);
-        if (!all_done)
+        done = hoomd::detail::WarpScan<bool, threads_per_particle>().Broadcast(done, 0);
+        if (!done)
             {
             // scan over flags
             unsigned char k(0), n(0);

                else
                    {
                    // we are past the end of the cell neighbors
                    this_thread_done = true;
                    }
                }
            if (! this_thread_done)
                {
                neigh_cell = d_cell_adj[cadji(cur_adj, my_cell)];
                neigh_size = d_cell_size[neigh_cell+igpu*ci.getNumElements()];
                }
            }
        // check for a neighbor if thread is still working
        if (!this_thread_done)
            {
            Scalar4 cur_xyzf;
            unsigned int j;

    // current index in cell
    int cur_offset = threadIdx.x % threads_per_particle;

    bool done = false;

    // total number of neighbors
    unsigned int nneigh = 0;

    while (! done)
        {
        // initialize with default
        unsigned int neighbor;
        unsigned char has_neighbor = 0;

        // advance neighbor cell
        while (cur_offset >= neigh_size && !done )
            {
            cur_offset -= neigh_size;
            cur_adj++;
