     
     // free GPU memory
     CUDA_SAFE_CALL(cudaFree(gd_pos));
+    CUDA_SAFE_CALL(cudaFree(gd_idxlist_coord));
+    CUDA_SAFE_CALL(cudaFree(gd_idxlist_coord_trans));
     CUDA_SAFE_CALL(cudaFreeArray(gd_idxlist_coord_array));
     CUDA_SAFE_CALL(cudaFreeArray(gd_idxlist_coord_trans_array));
     CUDA_SAFE_CALL(cudaFree(gd_bin_size));

    /*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>();*/
    
    printf("sorting bin_adj:\n");
    sort_bin_adj();

    cudaChannelFormatDesc channelDescFloat4 = cudaCreateChannelDesc<float4>();
    CUDA_SAFE_CALL(cudaMallocArray(&gd_idxlist_coord_trans_array, &channelDescFloat4,  Nbins, g_Nmax));
    CUDA_SAFE_CALL(cudaMallocArray(&gd_idxlist_coord_array, &channelDescFloat4, g_Nmax, Nbins));
    CUDA_SAFE_CALL(cudaMalloc((void**)&gd_idxlist_coord, Nbins * g_Nmax * sizeof(float4)));
    CUDA_SAFE_CALL(cudaMalloc((void**)&gd_idxlist_coord_trans, Nbins * g_Nmax * sizeof(float4)));
    
    gh_bin_size = (unsigned int *)malloc(Nbins * sizeof(unsigned int));
    CUDA_SAFE_CALL(cudaMalloc((void**)&gd_bin_size, Nbins * g_Nmax * sizeof(unsigned int)));
