     if (transpose_idxlist_coord)
         idxlist_coord_width = Mx*My*Mz;
     else
-        idxlist_coord_width = Nmax; 
+        idxlist_coord_width = Nmax;
+
+    unsigned int bin_adj_width = 0;
+    if (transpose_bin_adj)
+        bin_adj_width = Mx*My*Mz;
+    else
+        bin_adj_width = 27;
 
     // first, determine which bin this particle belongs to
     // MEM TRANSFER: 32 bytes

    
    // free GPU memory
    CUDA_SAFE_CALL(cudaFree(gd_pos));
    CUDA_SAFE_CALL(cudaFree(gd_idxlist_coord));
    CUDA_SAFE_CALL(cudaFree(gd_idxlist_coord_trans));
    CUDA_SAFE_CALL(cudaFreeArray(gd_idxlist_coord_array));

    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, 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, bin_size, d_bin_adj);
    }

// benchmark the device neighborlist
