         if (master && group == 0)
             s_still_searching = 0;
 
-        unsigned int tidx_1d = threadIdx.x+blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;
+        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))

                                                                 args.d_active_cell_accept,
                                                                 args.d_active_cell_move_type_translate,
                                                                 params,
                                                                 block_size);

    return cudaSuccess;
    }

        }

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