 unsigned int *gh_bin_size;  // number of particles in each bin
 
 unsigned int *gh_bin_adj;   // 27 x nbins adjacency array
+unsigned int *gd_bin_adj;   // 27 x nbins adjacency array
 cudaArray *gd_bin_adj_array;    // 27 x nbins adjancency array
 unsigned int *gh_bin_adj_trans; // nbins x 27 adjacency array
+unsigned int *gd_bin_adj_trans; // nbins x 27 adjacency array
 cudaArray *gd_bin_adj_trans_array;  // nbins x 27 adjancency array
 
 uint4 *gd_bin_coords;   // pre-calculated bin coordinates for each bin


#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, 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;

    if (transpose_idxlist_coord)
        idxlist_coord_width = Mx*My*Mz;
    else
        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
