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

//! Texture for reading particle charges
scalar_tex_t pdata_charge_tex;

//! Texture for reading neighbor list
texture<unsigned int, 1, cudaReadModeElementType> pdata_nlist_tex;

//! Kernel for calculating pair forces (shared memory version)
/*! This kernel is called to calculate the pair forces on all N particles. Actual evaluation of the potentials and

                }

            // 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;
            }
