     cell. Otherwise, extremely non-uniform cell lengths (i.e. avg 1, max 4) don't cause massive performance degradation.
 
     **Indexing**
-        - threadIdx.z indexes the current group in the block
+        - threadIdx.y indexes the current group in the block
         - threadIdx.x is the offset within the current group
         - blockIdx.x runs enough blocks so that all active cells are covered
 

            }

        n_groups = block_size / group_size / stride;
        shared_bytes = n_groups * (sizeof(unsigned int)*2 + sizeof(Scalar4) + sizeof(Scalar3)) +
                       block_size*(sizeof(unsigned int) + sizeof(unsigned int)) +
                       min_shared_bytes;
        }


        if (master && group == 0)
            s_still_searching = 0;

        unsigned int tidx_1d = threadIdx.x+blockDim.x*threadIdx.y;  // z component is for Shape parallelism

        // max_queue_size is always <= block size, so we just need an if here
        if (tidx_1d < min(s_queue_size, max_queue_size))
