     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
 

        }

    unsigned int n_groups = block_size / group_size / stride;
    unsigned int shared_bytes = n_groups * (sizeof(unsigned int)*2 + sizeof(Scalar4) + sizeof(Scalar3)) +
                                block_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;

            }

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

