     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;
+    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)

                shared_bytes += sizeof(Scalar)*block_size;

                gpu_compute_pair_forces_shared_kernel<evaluator, 0, 1>
                  <<<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);
                break;
                }
            case 1:

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

    if (idx >= const_storage_pair.N)
        return;

    // load in the length of the neighbor list (MEM_TRANSFER: 4 bytes)
    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);
