         }
 
     unsigned int n_groups = block_size / group_size / stride;
+    unsigned int max_queue_size = n_groups*group_size;
     unsigned int shared_bytes = n_groups * (sizeof(unsigned int)*2 + sizeof(Scalar4) + sizeof(Scalar3)) +
-                                block_size*(sizeof(unsigned int) + sizeof(unsigned int)) +
+                                max_queue_size*(sizeof(unsigned int) + sizeof(unsigned int)) +
                                 args.num_types * (sizeof(typename Shape::param_type) + 2*sizeof(Scalar)) +
                                 args.overlap_idx.getNumElements() * sizeof(unsigned int) +
                                 extra_bytes;

    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.x is the offset within the current group
        - blockIdx.x runs enough blocks so that all active cells are covered


        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))
