         ( _d_edges, cv::cuda::PtrStepSzb(_d_intermediate) );
 
     uint32_t edge_counter = 0;
-    POP_CUDA_MEMCPY_ASYNC( _d_edge_counter, &edge_count, sizeof(uint32_t), cudaMemcpyHostToDevice, _stream );
+    POP_CUDA_MEMCPY_ASYNC( _d_edge_counter, &edge_count, sizeof(uint32_t), cudaMemcpyHostToDevice, _stream, true );
 
     thinning_and_store
         <<<grid,block,0,_stream>>>
         ( cv::cuda::PtrStepSzb(_d_intermediate), _d_edges, _d_edge_counter, _edge_max, _d_edgelist );
 
-    POP_CUDA_MEMCPY_ASYNC( &edge_counter, _d_edge_counter, sizeof(uint32_t), cudaMemcpyDeviceToHost, _stream );
+    POP_CUDA_MEMCPY_ASYNC( &edge_counter, _d_edge_counter, sizeof(uint32_t), cudaMemcpyDeviceToHost, _stream, true );
 
     // Note: right here, Dynamic Parallelism would avoid blocking.
     cudaStreamSynchronize( _stream );


    {
        uint32_t dummy = 0;
        POP_CUDA_MEMCPY_ASYNC( _d_edge_counter, &dummy, sizeof(uint32_t), cudaMemcpyHostToDevice, _stream );
    }

    gradiant_descent
        <<<grid,block,0,_stream>>>
        ( _d_edgelist, edge_counter, _d_edgelist_2, _d_edge_counter, _d_edges, nmax, _d_dx, _d_dy, threshold, _d_out_points );

    POP_CUDA_MEMCPY_ASYNC( &edge_counter, _d_edge_counter, sizeof(uint32_t), cudaMemcpyDeviceToHost,_stream );

    // Note: right here, Dynamic Parallelism would avoid blocking.
    cudaStreamSynchronize( _stream );

    write_index += __popc( mask & ((1 << threadIdx.x) - 1) ); // find own write index

    if( keep && write_index < edgeMax ) {
        edgeCoords[write_index] = make_int2( idx, idy );
    }
}

