     pdata_charge_tex.filterMode = cudaFilterModePoint;
     cudaBindTexture(0, pdata_charge_tex, pair_args.d_charge, sizeof(Scalar) * pair_args.n_max);
     
+    // bind the head list texture
+//     pdata_head_list_tex.normalized = false;
+//     pdata_head_list_tex.filterMode = cudaFilterModePoint;
+//     cudaBindTexture(0, pdata_head_list_tex, pair_args.d_head_list, sizeof(unsigned int) * 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);

//! Texture for reading particle charges
scalar_tex_t pdata_charge_tex;

//! Texture for reading neighbor list
static texture<unsigned int> 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

            // read the current neighbor index (MEM TRANSFER: 4 bytes)
            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];

            // get the neighbor's position (MEM TRANSFER: 16 bytes)
            Scalar4 postypej = texFetchScalar4(d_pos, pdata_pos_tex, cur_j);
