     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, unsigned int *d_bin_adj)
-    {
-    const int block_size = 64;
-    
+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, unsigned int *d_bin_adj, unsigned int block_size)
+    {
     // setup the grid to run the kernel
     int nblocks = (int)ceil((double)N/ (double)block_size);
     

        
    // 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
    CUDA_SAFE_CALL(cudaMemcpy(gh_n_neigh, gd_n_neigh, g_N*sizeof(unsigned int), cudaMemcpyDeviceToHost));
   
    // verify results
    /*if (!verify())
        {
        printf("Invalid results in device bmark!\n");
        return 0.0f;
        }*/
        
    // benchmarks
    timeval start;
    cudaThreadSynchronize();
    gettimeofday(&start, NULL);
    
    unsigned int iters = 100;
    for (unsigned int i = 0; i < iters; i++)
        {
        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);
        }
        
    cudaThreadSynchronize();
    timeval end;
    gettimeofday(&end, NULL);
    float t = (end.tv_sec - start.tv_sec)*1000.0f + (end.tv_usec - start.tv_usec)/1000.0f;
    float avg_t = t/float(iters);
    
    printf("Device<%1d,%1d>          : ", transpose_idxlist_coord, transpose_bin_adj);
    printf("%f ms\n", avg_t);
    return avg_t;
    }

void bmark_grid()

float g_rcut;
const unsigned int g_Nmax = 64;    // Maximum number of particles each cell can hold
const float tweak_dist = 0.1f;
const unsigned int g_neigh_max = 137;

//*************** data structures
float4 *gh_pos, *gd_pos;            // particle positions
