         idx = blockIdx.x * (blockDim.x/tpp) + threadIdx.x/tpp;
         }
 
-    if (idx >= N)
+    if (idx >= const_storage_pair.N)
         return;
 
     // load in the length of the neighbor list (MEM_TRANSFER: 4 bytes)
-    unsigned int n_neigh = d_n_neigh[idx];
+    unsigned int n_neigh = const_storage_pair.d_n_neigh[idx];
 
     // read in the position of our particle.
     // (MEM TRANSFER: 16 bytes)
+    const Scalar4 *d_pos = const_storage_pair.d_pos;
     Scalar4 postypei = texFetchScalar4(d_pos, pdata_pos_tex, idx);
     Scalar3 posi = make_scalar3(postypei.x, postypei.y, postypei.z);
 
+    const Scalar *d_diameter = const_storage_pair.d_diameter;
     Scalar di;
     if (evaluator::needsDiameter())
         di = texFetchScalar(d_diameter, pdata_diam_tex, idx);
     else
         di += Scalar(1.0); // shutup compiler warning
+    const Scalar *d_charge = const_storage_pair.d_charge;
     Scalar qi;
     if (evaluator::needsCharge())
         qi = texFetchScalar(d_charge, pdata_charge_tex, idx);

    //! Construct a pair_args_t
    pair_args_t(Scalar4 *_d_force,
              Scalar *_d_virial,
              const unsigned int _virial_pitch,
              const unsigned int _N,
              const unsigned int _n_max,
              const Scalar4 *_d_pos,
              const Scalar *_d_diameter,
              const Scalar *_d_charge,
              const BoxDim& _box,
              const unsigned int *_d_n_neigh,
              const unsigned int *_d_nlist,
              const Index2D& _nli,
              const Scalar *_d_rcutsq,
              const Scalar *_d_ronsq,
              const unsigned int _ntypes,
              const unsigned int _block_size,
              const unsigned int _shift_mode,
              const unsigned int _compute_virial,
              const unsigned int _threads_per_particle)
                : d_force(_d_force),
                  d_virial(_d_virial),
                  virial_pitch(_virial_pitch),

    Scalar virialyz = Scalar(0.0);
    Scalar virialzz = Scalar(0.0);

    // load parameters
    Index2D nli(const_storage_pair.nli_w,const_storage_pair.nli_h);
    const unsigned int *d_nlist = const_storage_pair.d_nlist;

    unsigned int cur_j = 0;
    unsigned int next_j = threadIdx.x%tpp < n_neigh ?  d_nlist[nli(idx, threadIdx.x%tpp)] : 0;
 
    // loop over neighbors
    // on pre Fermi hardware, there is a bug that causes rare and random ULFs when simply looping over n_neigh
    // the workaround (activated via the template paramter) is to loop over nlist.height and put an if (i < n_neigh)
