                                                  unsigned int window_size,
                                                  unsigned int thread_mask,
                                                  unsigned int n_bodies_per_block,
-                                                 bool zero_force)
+                                                 bool zero_force,
+                                                 unsigned int first_body,
+                                                 unsigned int nwork)
     {
     // determine which body (0 ... n_bodies_per_block-1) this thread is working on
     // assign threads 0, 1, 2, ... to body 0, n, n+1, n+2, ... to body 1, and so on.

            window_size,
            thread_mask,
            n_bodies_per_block,
            zero_force);
        }
    return cudaSuccess;
    }

        {
        // thread 0 for this body reads in the body id and orientation and stores them in shared memory
        int group_idx = blockIdx.x*n_bodies_per_block + m;
        if (group_idx < nwork)
            {
            central_idx[m] = d_rigid_center[group_idx + first_body];
            mol_idx[m] = d_molecule_idx[central_idx[m]];

            if (d_tag[central_idx[m]] != d_body[central_idx[m]])
