 
 #define EMPTY_BIN 0xffffffff
 
-template<bool transpose_idxlist_coord, bool transpose_bin_adj> __global__ void gpu_compute_nlist_binned_kernel(unsigned int *d_nlist, unsigned int *d_n_neigh, unsigned int nlist_pitch, unsigned int neigh_max, float4 *d_pos, unsigned int local_num, float Lx, float Ly, float Lz, float Lxinv, float Lyinv, float Lzinv, float r_maxsq, unsigned int actual_Nmax, float scalex, float scaley, float scalez, unsigned int Mx, unsigned int My, unsigned int Mz)
+template<bool transpose_idxlist_coord, bool transpose_bin_adj> __global__ void gpu_compute_nlist_binned_kernel(unsigned int *d_nlist, unsigned int *d_n_neigh, unsigned int nlist_pitch, unsigned int neigh_max, float4 *d_pos, float4 *d_idxlist_coord, unsigned int Nmax, unsigned int local_num, float Lx, float Ly, float Lz, float Lxinv, float Lyinv, float Lzinv, float r_maxsq, unsigned int actual_Nmax, float scalex, float scaley, float scalez, unsigned int Mx, unsigned int My, unsigned int Mz)
     {
     // each thread is going to compute the neighbor list for a single particle
     int idx = blockDim.x * blockIdx.x + threadIdx.x;

        bin_adj_array = gd_bin_adj_array;
        
    // warm up
    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, g_N, g_Lx, g_Ly, g_Lz, g_Mx, g_My, g_Mz, g_Nmax);
    CUT_CHECK_ERROR("kernel failed");
    
    // copy results back

    /*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();
    
