 //! 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
 texture<unsigned int, 1, cudaReadModeElementType> pdata_nlist_tex;
 

    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;
    

    pdata_charge_tex.filterMode = cudaFilterModePoint;
    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;
