 /*! 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 d_force Device memory to write computed forces
-    \param d_virial Device memory to write computed virials
-    \param virial_pitch pitch of 2D virial array
-    \param N number of particles in system
-    \param d_pos particle positions
-    \param d_diameter particle diameters
-    \param d_charge particle charges
     \param box Box dimensions used to implement periodic boundary conditions
-    \param d_n_neigh Device memory array listing the number of neighbors for each particle
-    \param d_nlist Device memory array containing the neighbor list contents
-    \param nli Indexer for indexing \a d_nlist
     \param d_params Parameters for the potential, stored per type pair
-    \param d_rcutsq rcut squared, stored per type pair
-    \param d_ronsq ron squared, stored per type pair
-    \param ntypes Number of types in the simulation
 
     \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

        idx = blockIdx.x * (blockDim.x/tpp) + threadIdx.x/tpp;
        }

    if (idx >= N)
        return;

    // load in the length of the neighbor list (MEM_TRANSFER: 4 bytes)
    unsigned int n_neigh = d_n_neigh[idx];

    // read in the position of our particle.
    // (MEM TRANSFER: 16 bytes)
    Scalar4 postypei = texFetchScalar4(d_pos, pdata_pos_tex, idx);
    Scalar3 posi = make_scalar3(postypei.x, postypei.y, postypei.z);

    Scalar di;
    if (evaluator::needsDiameter())
        di = texFetchScalar(d_diameter, pdata_diam_tex, idx);
    else
        di += Scalar(1.0); // shutup compiler warning
    Scalar qi;
    if (evaluator::needsCharge())
        qi = texFetchScalar(d_charge, pdata_charge_tex, idx);

        // if we are the first thread in the cta, write out virial to global mem
        if (threadIdx.x %tpp == 0)
            {
            Scalar *d_virial = const_storage_pair.d_virial;
            unsigned int virial_pitch = const_storage_pair.virial_pitch;
            d_virial[0*virial_pitch+idx] = virialxx;
            d_virial[1*virial_pitch+idx] = virialxy;
            d_virial[2*virial_pitch+idx] = virialxz;
