 //! Texture for reading particle charges
 scalar_tex_t pdata_charge_tex;
 
+//! Texture for reading head list
+// texture<unsigned int, 1, cudaReadModeElementType> pdata_head_list_tex;
+
 //! Texture for reading neighbor list
-static texture<unsigned int> pdata_nlist_tex;
+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

    pdata_charge_tex.filterMode = cudaFilterModePoint;
    cudaBindTexture(0, pdata_charge_tex, pair_args.d_charge, sizeof(Scalar) * pair_args.n_max);
    
    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);

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

//     #if __CUDA_ARCH__ >= 350
//     unsigned int my_head = __ldg(d_head_list + idx);
//     #else
//     unsigned int my_head = tex1Dfetch(pdata_head_list_tex, idx);
//     #endif
    unsigned int my_head = d_head_list[idx];
    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 ?
//         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)
        {
