     float scalez = 1.0f / binz;
     
     // run the kernel
-    gpu_compute_nlist_binned_kernel<transpose_idxlist_coord, transpose_bin_adj><<<grid,threads>>>(nlist, n_neigh, nlist_pitch, neigh_max, pos, N, Lx, Ly, Lz, 1.0f/Lx, 1.0f/Ly, 1.0f/Lz, r_cut_sq, g_actual_Nmax, scalex, scaley, scalez, Mx, My, Mz);
+    gpu_compute_nlist_binned_kernel<transpose_idxlist_coord, transpose_bin_adj><<<grid,threads>>>(nlist, n_neigh, nlist_pitch, neigh_max, pos, d_idxlist_coord, Nmax, N, Lx, Ly, Lz, 1.0f/Lx, 1.0f/Ly, 1.0f/Lz, r_cut_sq, g_actual_Nmax, scalex, scaley, scalez, Mx, My, Mz);
     }
 
 // benchmark the device neighborlist
 template<bool transpose_idxlist_coord, bool transpose_bin_adj> void bmark_device_nlist()
     {
     cudaArray *idxlist_coord_array;
+    float4 *d_idxlist_coord;
     if (transpose_idxlist_coord)
+        {
         idxlist_coord_array = gd_idxlist_coord_trans_array;
+        d_idxlist_coord = gd_idxlist_coord_trans;
+        }
     else
+        {
         idxlist_coord_array = gd_idxlist_coord_array;
+        d_idxlist_coord = gd_idxlist_coord;
+        }
         
     cudaArray *bin_adj_array;
     if (transpose_bin_adj)

            cout << setw(4) << gh_bin_adj[27*cur_bin + i] << " ";
        cout << endl;
        }
    cout << endl;
    
    // sort the bin_adj lists
    for (unsigned int cur_bin = 0; cur_bin < g_Mx*g_My*g_Mz; cur_bin++)

    /*bmark_host_nlist<false, false>();
    bmark_host_nlist<false, true>();
    bmark_host_nlist<true, false>();
    bmark_host_nlist<true, true>();*/
    
    bmark_device_nlist<false, false>();
    bmark_device_nlist<false, true>();
    bmark_device_nlist<true, false>();
    bmark_device_nlist<true, true>();
    
    free_data();
    
