             cur_j = next_j;
             if (neigh_idx+tpp < n_neigh)
                 {
-                #if __CUDA_ARCH__ >= 350
-                next_j = __ldg(d_nlist + my_head + neigh_idx + tpp);
-                #else
-                next_j = tex1Dfetch(pdata_nlist_tex, my_head + neigh_idx + tpp);
-                #endif
-                }
-//                 next_j = d_nlist[my_head + neigh_idx+tpp];
+                next_j = texFetchUint(d_nlist, pair_nlist_tex, my_head + neigh_idx + tpp);
+                }
 
             // get the neighbor's position (MEM TRANSFER: 16 bytes)
             Scalar4 postypej = texFetchScalar4(d_pos, pdata_pos_tex, cur_j);

                }

            // add up the force vector components (FLOPS: 7)
            #if (__CUDA_ARCH__ >= 200)
            force.x += dx.x * force_divr;
            force.y += dx.y * force_divr;
            force.z += dx.z * force_divr;
            #else
            // fmad causes momentum drift here, prevent it from being used
            force.x += __fmul_rn(dx.x, force_divr);
            force.y += __fmul_rn(dx.y, force_divr);
            force.z += __fmul_rn(dx.z, force_divr);
            #endif

            force.w += pair_eng;
            }

    const unsigned int *d_head_list;//!< Head list indexes 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 size_neigh_list; //!< Size of the neighbor list for texture binding
    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
