                 shared_bytes += sizeof(Scalar)*block_size;
 
                 gpu_compute_pair_forces_shared_kernel<evaluator, 0, 0>
-                  <<<grid, block_size, shared_bytes>>>(pair_args.d_force, pair_args.d_virial,
-                  pair_args.virial_pitch, pair_args.N, pair_args.d_pos, pair_args.d_diameter,
-                  pair_args.d_charge, pair_args.box, pair_args.d_n_neigh, pair_args.d_nlist,
-                  pair_args.nli, d_params, pair_args.d_rcutsq, pair_args.d_ronsq, pair_args.ntypes,
-                  tpp);
+                    <<<grid, block_size, shared_bytes>>>(pair_args.box,d_params);
                 break;
                 }
             case 1:


    Scalar4 *d_force;                //!< Force to write out
    Scalar *d_virial;                //!< Virial to write out
    const unsigned int virial_pitch; //!< The pitch of the 2D array of virial matrix elements
    const unsigned int N;           //!< number of particles
    const unsigned int n_max;       //!< Max size of pdata arrays
    const Scalar4 *d_pos;           //!< particle positions
    const Scalar *d_diameter;       //!< particle diameters
    const Scalar *d_charge;         //!< particle charges
    const BoxDim& box;         //!< Simulation box in GPU format
    const unsigned int *d_n_neigh;  //!< Device array listing the number of neighbors on each particle
    const unsigned int *d_nlist;    //!< Device array listing the neighbors of each particle
    const Index2D& nli;             //!< Indexer for accessing d_nlist
    const Scalar *d_rcutsq;          //!< Device array listing r_cut squared per particle type pair
    const Scalar *d_ronsq;           //!< Device array listing r_on squared per particle type pair
    const unsigned int ntypes;      //!< Number of particle types in the simulation
    const unsigned int block_size;  //!< Block size to execute
    const unsigned int shift_mode;  //!< The potential energy shift mode
    const unsigned int compute_virial;  //!< Flag to indicate if virials should be computed
    const unsigned int threads_per_particle; //!< Number of threads per particle (maximum: 1 warp)
    };

#ifdef NVCC
//! Texture for reading particle positions

/*! This kernel is called to calculate the pair forces on all N particles. Actual evaluation of the potentials and
    forces for each pair is handled via the template class \a evaluator.

    \param box Box dimensions used to implement periodic boundary conditions
    \param d_params Parameters for the potential, stored per type pair

    \a d_params, \a d_rcutsq, and \a d_ronsq must be indexed with an Index2DUpperTriangler(typei, typej) to access the
    unique value for that type pair. These values are all cached into shared memory for quick access, so a dynamic
