     extern __shared__ char s_data[];
     typename Shape::param_type *s_params = (typename Shape::param_type *)(&s_data[0]);
     unsigned int *s_check_overlaps = (unsigned int *) (s_params + num_types);
-
     unsigned int ntyppairs = overlap_idx.getNumElements();
+    unsigned int *s_overlap = (unsigned int *)(&s_check_overlaps[ntyppairs]);
+
     // copy over parameters one int per thread for fast loads
         {
         unsigned int tidx = threadIdx.x+blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z;

        grid.x = 65535;
        }

    unsigned int shared_bytes = args.num_types * sizeof(typename Shape::param_type) + n_groups*sizeof(unsigned int)
        + args.overlap_idx.getNumElements()*sizeof(unsigned int);

    gpu_hpmc_free_volume_kernel<Shape><<<grid, threads, shared_bytes>>>(
                                                     args.n_sample,

            }
        }

    __syncthreads();

    // initialize extra shared mem
    char *s_extra = (char *)(s_overlap + n_groups);

    for (unsigned int cur_type = 0; cur_type < num_types; ++cur_type)
        s_params[cur_type].load_shared(s_extra, true);

    __shared__ unsigned int s_n_overlap;

    if (master)
