         
     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));
-    
+    CUDA_SAFE_CALL(cudaMemcpy(gd_bin_adj, gh_bin_adj, sizeof(unsigned int)*g_Mx*g_My*g_Mz*27, cudaMemcpyHostToDevice));
+    CUDA_SAFE_CALL(cudaMemcpy(gd_bin_adj_trans, 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);

    d_n_neigh[my_pidx] = n_neigh;
    }

template<bool transpose_idxlist_coord, bool transpose_bin_adj> void gpu_compute_nlist_binned(unsigned int *nlist, unsigned int *n_neigh, unsigned int nlist_pitch, float r_cut_sq, unsigned int neigh_max, cudaArray *idxlist_coord, unsigned int *bin_size, uint4 *bin_coords, cudaArray *bin_adj, float4 *pos, float4 *d_idxlist_coord, unsigned int N, float Lx, float Ly, float Lz, unsigned int Mx, unsigned int My, unsigned int Mz, unsigned int Nmax)
    {
    const int block_size = 64;
    

        }
        
    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
