     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);
+    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

        
    CUDA_SAFE_CALL(cudaMemcpyToArray(gd_bin_adj_array, 0, 0, gh_bin_adj, sizeof(unsigned int)*g_Mx*g_My*g_Mz*27, cudaMemcpyHostToDevice));
    CUDA_SAFE_CALL(cudaMemcpyToArray(gd_bin_adj_trans_array, 0, 0, gh_bin_adj_trans, sizeof(unsigned int)*g_Mx*g_My*g_Mz*27, cudaMemcpyHostToDevice));
    
    memset(gh_nlist, 0, sizeof(unsigned int)*g_nlist_pitch*g_neigh_max);
    memset(gh_n_neigh, 0, sizeof(unsigned int)*g_N);
    memset(gh_nlist_ref, 0, sizeof(unsigned int)*g_nlist_pitch*g_neigh_max);

        }
        
    cudaArray *bin_adj_array;
    unsigned int *d_bin_adj;
    if (transpose_bin_adj)
        {
        bin_adj_array = gd_bin_adj_trans_array;
        d_bin_adj = gd_bin_adj_trans;
        }
    else
        {
        bin_adj_array = gd_bin_adj_array;
        d_bin_adj = gd_bin_adj;
        }
        
    // warm up
    CUDA_SAFE_CALL(cudaFuncSetCacheConfig(gpu_compute_nlist_binned_kernel<transpose_idxlist_coord, transpose_bin_adj>, cudaFuncCachePreferL1));
    gpu_compute_nlist_binned<transpose_idxlist_coord, transpose_bin_adj>(gd_nlist, gd_n_neigh, g_nlist_pitch, g_rcut*g_rcut, g_neigh_max, idxlist_coord_array, gd_bin_size, gd_bin_coords, bin_adj_array, gd_pos, d_idxlist_coord, g_N, g_Lx, g_Ly, g_Lz, g_Mx, g_My, g_Mz, g_Nmax, d_bin_adj);
    CUT_CHECK_ERROR("kernel failed");
    
    // copy results back
