     unsigned int cur_j = 0;
     
     // first read only neighborlist fetch
-    #if __CUDA_ARCH__ >= 350
-    unsigned int next_j = threadIdx.x%tpp < n_neigh ?
-        __ldg(d_nlist + my_head + threadIdx.x%tpp) : 0;
-    #else
-    unsigned int next_j = threadIdx.x%tpp < n_neigh ?
-        tex1Dfetch(pdata_nlist_tex, my_head + threadIdx.x%tpp) : 0;
-    #endif
+    unsigned int next_j = threadIdx.x%tpp < n_neigh ? texFetchUint(d_nlist, pair_nlist_tex, my_head + threadIdx.x%tpp) : 0;
     
-//     unsigned int next_j = threadIdx.x%tpp < n_neigh ?
-//         d_nlist[my_head + threadIdx.x%tpp] : 0;
     // loop over neighbors
     for (int neigh_idx = threadIdx.x%tpp; neigh_idx < n_neigh; neigh_idx+=tpp)
         {

    cudaBindTexture(0, pdata_charge_tex, pair_args.d_charge, sizeof(Scalar) * pair_args.n_max);
    
    // bind the neighborlist texture
    pdata_nlist_tex.normalized = false;
    pdata_nlist_tex.filterMode = cudaFilterModePoint;
    cudaBindTexture(0, pdata_nlist_tex, pair_args.d_nlist, sizeof(unsigned int) * pair_args.size_neigh_list);
    }

//! Kernel driver that computes lj forces on the GPU for LJForceComputeGPU

                }

            // add up the force vector components (FLOPS: 7)
            force.x += dx.x * force_divr;
            force.y += dx.y * force_divr;
            force.z += dx.z * force_divr;

            force.w += pair_eng;
            }
