Repository: NVIDIA-Genomics-Research/GenomeWorks (language: python)

Issue:
[common] Rename device_copy_n to reflect non-blocking behavior
The version of `device_copy_n` which takes a CUDA stream argument
https://github.com/clara-parabricks/GenomeWorks/blob/74e6424c156a7ee15b4137d4788a4257ee6482c4/common/base/include/claraparabricks/genomeworks/utils/cudautils.hpp#L138
should be renamed to `device_copy_n_async` to make the non-blocking behavior more obvious.


The blocking three-argument version of `device_copy_n` should remain as is.


Discussion / hints:
Furthermore, `device_copy_n` and `device_copy_n_async` should be moved one namespace up - from `claraparabricks::genomeworks::cudautils` to `claraparabricks::genomeworks`.

Fix patch:
diff --git a/common/base/include/claraparabricks/genomeworks/utils/cudautils.hpp b/common/base/include/claraparabricks/genomeworks/utils/cudautils.hpp
index e826bcb0d..82b9d092c 100644
--- a/common/base/include/claraparabricks/genomeworks/utils/cudautils.hpp
+++ b/common/base/include/claraparabricks/genomeworks/utils/cudautils.hpp
@@ -136,7 +136,7 @@ void set_device_value(Type* dst, const Type& src)
 
 /// Copies elements from the range [src, src + n) to the range [dst, dst + n) asynchronously.
 template <typename Type>
-void device_copy_n(const Type* src, size_t n, Type* dst, cudaStream_t stream)
+void device_copy_n_async(const Type* src, size_t n, Type* dst, cudaStream_t stream)
 {
     GW_CU_CHECK_ERR(cudaMemcpyAsync(dst, src, n * sizeof(Type), cudaMemcpyDefault, stream));
 }
diff --git a/cudaaligner/src/aligner_global_myers_banded.cpp b/cudaaligner/src/aligner_global_myers_banded.cpp
index 494152fc0..8eb6eff46 100644
--- a/cudaaligner/src/aligner_global_myers_banded.cpp
+++ b/cudaaligner/src/aligner_global_myers_banded.cpp
@@ -240,7 +240,7 @@ StatusType AlignerGlobalMyersBanded::add_alignment(const char* query, int32_t qu
 
 StatusType AlignerGlobalMyersBanded::align_all()
 {
-    using cudautils::device_copy_n;
+    using cudautils::device_copy_n_async;
     const auto n_alignments = get_size(alignments_);
     if (n_alignments == 0)
         return StatusType::success;
@@ -278,9 +278,9 @@ StatusType AlignerGlobalMyersBanded::align_all()
     {
         result_lengths_d.clear_and_resize(n_alignments);
     }
-    device_copy_n(seq_h.data(), seq_starts_h.back(), seq_d.data(), stream_);
-    device_copy_n(seq_starts_h.data(), 2 * n_alignments + 1, seq_starts_d.data(), stream_);
-    device_copy_n(result_starts_h.data(), n_alignments + 1, result_starts_d.data(), stream_);
+    device_copy_n_async(seq_h.data(), seq_starts_h.back(), seq_d.data(), stream_);
+    device_copy_n_async(seq_starts_h.data(), 2 * n_alignments + 1, seq_starts_d.data(), stream_);
+    device_copy_n_async(result_starts_h.data(), n_alignments + 1, result_starts_d.data(), stream_);
 
     myers_banded_gpu(results_d.data(), result_lengths_d.data(), result_starts_d.data(),
                      seq_d.data(), seq_starts_d.data(), n_alignments, max_bandwidth_,
@@ -290,8 +290,8 @@ StatusType AlignerGlobalMyersBanded::align_all()
     result_lengths_h.clear();
     result_lengths_h.resize(n_alignments);
 
-    device_copy_n(results_d.data(), result_starts_h.back(), results_h.data(), stream_);
-    device_copy_n(result_lengths_d.data(), n_alignments, result_lengths_h.data(), stream_);
+    device_copy_n_async(results_d.data(), result_starts_h.back(), results_h.data(), stream_);
+    device_copy_n_async(result_lengths_d.data(), n_alignments, result_lengths_h.data(), stream_);
 
     return StatusType::success;
 }
diff --git a/cudaaligner/src/batched_device_matrices.cuh b/cudaaligner/src/batched_device_matrices.cuh
index e0758f1ce..773fa9374 100644
--- a/cudaaligner/src/batched_device_matrices.cuh
+++ b/cudaaligner/src/batched_device_matrices.cuh
@@ -215,8 +215,8 @@ public:
         offsets_.clear_and_resize(offsets_host_.size());
         dev_interface_host_.clear();
         dev_interface_host_.emplace_back(storage_.data(), offsets_.data(), get_size<int32_t>(offsets_host_) - 1);
-        cudautils::device_copy_n(offsets_host_.data(), offsets_host_.size(), offsets_.data(), stream);
-        cudautils::device_copy_n(dev_interface_host_.data(), 1, dev_.data(), stream);
+        cudautils::device_copy_n_async(offsets_host_.data(), offsets_host_.size(), offsets_.data(), stream);
+        cudautils::device_copy_n_async(dev_interface_host_.data(), 1, dev_.data(), stream);
     }
 
     matrix<T> get_matrix(int32_t id, int32_t n_rows, int32_t n_cols, cudaStream_t stream)
@@ -231,7 +231,7 @@ public:
         matrix<T> m(n_rows, n_cols);
         if (n_rows * n_cols > offsets_host_[id + 1] - offsets_host_[id])
             throw std::runtime_error("Requested matrix size is larger than allocated memory on device.");
-        cudautils::device_copy_n(storage_.data() + offsets_host_[id], n_rows * n_cols, m.data(), stream);
+        cudautils::device_copy_n_async(storage_.data() + offsets_host_[id], n_rows * n_cols, m.data(), stream);
         GW_CU_CHECK_ERR(cudaStreamSynchronize(stream));
         return m;
     }
diff --git a/cudaextender/samples/sample_cudaextender.cpp b/cudaextender/samples/sample_cudaextender.cpp
index 53d4bd8e4..69fb29d96 100644
--- a/cudaextender/samples/sample_cudaextender.cpp
+++ b/cudaextender/samples/sample_cudaextender.cpp
@@ -154,9 +154,9 @@ int main(int argc, char* argv[])
         device_buffer<int32_t> d_num_ssp(1, allocator, stream0.get());
 
         // Async Memcopy all the input values to device
-        device_copy_n(h_encoded_query.data(), get_size(query_sequence), d_query.data(), stream0.get());
-        device_copy_n(h_encoded_target.data(), get_size(target_sequence), d_target.data(), stream0.get());
-        device_copy_n(h_seed_pairs.data(), get_size(h_seed_pairs), d_seed_pairs.data(), stream0.get());
+        device_copy_n_async(h_encoded_query.data(), get_size(query_sequence), d_query.data(), stream0.get());
+        device_copy_n_async(h_encoded_target.data(), get_size(target_sequence), d_target.data(), stream0.get());
+        device_copy_n_async(h_seed_pairs.data(), get_size(h_seed_pairs), d_seed_pairs.data(), stream0.get());
 
         // Launch the ungapped extender device pointer function
         ungapped_extender->extend_async(d_query.data(),
@@ -174,7 +174,7 @@ int main(int argc, char* argv[])
         const int32_t h_num_ssp = cudautils::get_value_from_device(d_num_ssp.data(), stream0.get());
         h_ssp.resize(h_num_ssp);
         // Copy data asynchronously
-        device_copy_n(d_ssp.data(), h_num_ssp, h_ssp.data(), stream0.get());
+        device_copy_n_async(d_ssp.data(), h_num_ssp, h_ssp.data(), stream0.get());
         cudaStreamSynchronize(stream0.get());
     }
     std::cerr << "Number of Scored Segment Pairs found: " << get_size(h_ssp) << std::endl;
diff --git a/cudaextender/src/ungapped_xdrop.cu b/cudaextender/src/ungapped_xdrop.cu
index 6aa2b438f..a06af704a 100644
--- a/cudaextender/src/ungapped_xdrop.cu
+++ b/cudaextender/src/ungapped_xdrop.cu
@@ -94,7 +94,7 @@ UngappedXDrop::UngappedXDrop(const int32_t* h_score_mat, const int32_t score_mat
     d_done_             = device_buffer<int32_t>(batch_max_ungapped_extensions_, allocator_, stream_);
     d_tmp_ssp_          = device_buffer<ScoredSegmentPair>(batch_max_ungapped_extensions_, allocator_, stream_);
     d_temp_storage_cub_ = device_buffer<int8_t>(cub_storage_bytes, allocator_, stream_);
-    device_copy_n(h_score_mat_.data(), score_mat_dim_, d_score_mat_.data(), stream_);
+    device_copy_n_async(h_score_mat_.data(), score_mat_dim_, d_score_mat_.data(), stream_);
 }
 
 StatusType UngappedXDrop::extend_async(const int8_t* d_query, const int32_t query_length,
@@ -196,9 +196,9 @@ StatusType UngappedXDrop::extend_async(const int8_t* h_query, const int32_t quer
     d_num_ssp_ = device_buffer<int32_t>(1, allocator_, stream_);
 
     // Async memcopy all the input values to device
-    device_copy_n(h_query, query_length, d_query_.data(), stream_);
-    device_copy_n(h_target, target_length, d_target_.data(), stream_);
-    device_copy_n(h_seed_pairs.data(), h_seed_pairs.size(), d_seed_pairs_.data(), stream_);
+    device_copy_n_async(h_query, query_length, d_query_.data(), stream_);
+    device_copy_n_async(h_target, target_length, d_target_.data(), stream_);
+    device_copy_n_async(h_seed_pairs.data(), h_seed_pairs.size(), d_seed_pairs_.data(), stream_);
 
     // Launch the ungapped extender device function
     return extend_async(d_query_.data(), query_length,
@@ -216,7 +216,7 @@ StatusType UngappedXDrop::sync()
         if (h_num_ssp > 0)
         {
             h_ssp_.resize(h_num_ssp);
-            device_copy_n(d_ssp_.data(), h_num_ssp, h_ssp_.data(), stream_);
+            device_copy_n_async(d_ssp_.data(), h_num_ssp, h_ssp_.data(), stream_);
             GW_CU_CHECK_ERR(cudaStreamSynchronize(stream_));
         }
         return StatusType::success;
diff --git a/cudamapper/src/index_gpu.cuh b/cudamapper/src/index_gpu.cuh
index 6e17a1927..20f964cad 100644
--- a/cudamapper/src/index_gpu.cuh
+++ b/cudamapper/src/index_gpu.cuh
@@ -645,12 +645,12 @@ IndexGPU<SketchElementImpl>::IndexGPU(DefaultDeviceAllocator allocator,
     number_of_reads_                     = index_host_copy->number_of_reads();
     number_of_basepairs_in_longest_read_ = index_host_copy->number_of_basepairs_in_longest_read();
 
-    cudautils::device_copy_n(index_host_copy->representations().data, index_host_copy->representations().size, representations_d_.data(), cuda_stream);
-    cudautils::device_copy_n(index_host_copy->read_ids().data, index_host_copy->read_ids().size, read_ids_d_.data(), cuda_stream);
-    cudautils::device_copy_n(index_host_copy->positions_in_reads().data, index_host_copy->positions_in_reads().size, positions_in_reads_d_.data(), cuda_stream);
-    cudautils::device_copy_n(index_host_copy->directions_of_reads().data, index_host_copy->directions_of_reads().size, directions_of_reads_d_.data(), cuda_stream);
-    cudautils::device_copy_n(index_host_copy->unique_representations().data, index_host_copy->unique_representations().size, unique_representations_d_.data(), cuda_stream);
-    cudautils::device_copy_n(index_host_copy->first_occurrence_of_representations().data, index_host_copy->first_occurrence_of_representations().size, first_occurrence_of_representations_d_.data(), cuda_stream);
+    cudautils::device_copy_n_async(index_host_copy->representations().data, index_host_copy->representations().size, representations_d_.data(), cuda_stream);
+    cudautils::device_copy_n_async(index_host_copy->read_ids().data, index_host_copy->read_ids().size, read_ids_d_.data(), cuda_stream);
+    cudautils::device_copy_n_async(index_host_copy->positions_in_reads().data, index_host_copy->positions_in_reads().size, positions_in_reads_d_.data(), cuda_stream);
+    cudautils::device_copy_n_async(index_host_copy->directions_of_reads().data, index_host_copy->directions_of_reads().size, directions_of_reads_d_.data(), cuda_stream);
+    cudautils::device_copy_n_async(index_host_copy->unique_representations().data, index_host_copy->unique_representations().size, unique_representations_d_.data(), cuda_stream);
+    cudautils::device_copy_n_async(index_host_copy->first_occurrence_of_representations().data, index_host_copy->first_occurrence_of_representations().size, first_occurrence_of_representations_d_.data(), cuda_stream);
 }
 
 template <typename SketchElementImpl>
@@ -881,15 +881,15 @@ void IndexGPU<SketchElementImpl>::generate_index(const io::FastaParser& parser,
 
     {
         GW_NVTX_RANGE(profiler, "IndexGPU::generate_index::move_basepairs_to_gpu");
-        cudautils::device_copy_n(read_id_to_basepairs_section_h.data(),
-                                 read_id_to_basepairs_section_h.size(),
-                                 read_id_to_basepairs_section_d.data(),
-                                 cuda_stream); // H2D
-
-        cudautils::device_copy_n(merged_basepairs_h.data(),
-                                 merged_basepairs_h.size(),
-                                 merged_basepairs_d.data(),
-                                 cuda_stream); // H2D
+        cudautils::device_copy_n_async(read_id_to_basepairs_section_h.data(),
+                                       read_id_to_basepairs_section_h.size(),
+                                       read_id_to_basepairs_section_d.data(),
+                                       cuda_stream); // H2D
+
+        cudautils::device_copy_n_async(merged_basepairs_h.data(),
+                                       merged_basepairs_h.size(),
+                                       merged_basepairs_d.data(),
+                                       cuda_stream); // H2D
     }
 
     // sketch elements get generated here
diff --git a/cudamapper/src/index_host_copy.cu b/cudamapper/src/index_host_copy.cu
index 5a83e50dd..e4b2b4871 100644
--- a/cudamapper/src/index_host_copy.cu
+++ b/cudamapper/src/index_host_copy.cu
@@ -81,35 +81,35 @@ IndexHostCopy::IndexHostCopy(const Index& index,
     // register pinned memory, memory gets unpinned in finish_copying()
     memory_pinner_.register_pinned_memory();
 
-    cudautils::device_copy_n(index.representations().data(),
-                             index.representations().size(),
-                             representations_.data,
-                             cuda_stream_);
-
-    cudautils::device_copy_n(index.read_ids().data(),
-                             index.read_ids().size(),
-                             read_ids_.data,
-                             cuda_stream_);
-
-    cudautils::device_copy_n(index.positions_in_reads().data(),
-                             index.positions_in_reads().size(),
-                             positions_in_reads_.data,
-                             cuda_stream_);
-
-    cudautils::device_copy_n(index.directions_of_reads().data(),
-                             index.directions_of_reads().size(),
-                             directions_of_reads_.data,
-                             cuda_stream_);
-
-    cudautils::device_copy_n(index.unique_representations().data(),
-                             index.unique_representations().size(),
-                             unique_representations_.data,
-                             cuda_stream_);
-
-    cudautils::device_copy_n(index.first_occurrence_of_representations().data(),
-                             index.first_occurrence_of_representations().size(),
-                             first_occurrence_of_representations_.data,
-                             cuda_stream_);
+    cudautils::device_copy_n_async(index.representations().data(),
+                                   index.representations().size(),
+                                   representations_.data,
+                                   cuda_stream_);
+
+    cudautils::device_copy_n_async(index.read_ids().data(),
+                                   index.read_ids().size(),
+                                   read_ids_.data,
+                                   cuda_stream_);
+
+    cudautils::device_copy_n_async(index.positions_in_reads().data(),
+                                   index.positions_in_reads().size(),
+                                   positions_in_reads_.data,
+                                   cuda_stream_);
+
+    cudautils::device_copy_n_async(index.directions_of_reads().data(),
+                                   index.directions_of_reads().size(),
+                                   directions_of_reads_.data,
+                                   cuda_stream_);
+
+    cudautils::device_copy_n_async(index.unique_representations().data(),
+                                   index.unique_representations().size(),
+                                   unique_representations_.data,
+                                   cuda_stream_);
+
+    cudautils::device_copy_n_async(index.first_occurrence_of_representations().data(),
+                                   index.first_occurrence_of_representations().size(),
+                                   first_occurrence_of_representations_.data,
+                                   cuda_stream_);
 
     number_of_reads_                     = index.number_of_reads();
     number_of_basepairs_in_longest_read_ = index.number_of_basepairs_in_longest_read();
diff --git a/cudamapper/src/minimizer.cu b/cudamapper/src/minimizer.cu
index 6ee7286be..ada31d851 100644
--- a/cudamapper/src/minimizer.cu
+++ b/cudamapper/src/minimizer.cu
@@ -899,10 +899,10 @@ Minimizer::GeneratedSketchElements Minimizer::generate_sketch_elements(DefaultDe
     }
 
     device_buffer<decltype(read_id_to_windows_section_h)::value_type> read_id_to_windows_section_d(read_id_to_windows_section_h.size(), allocator, cuda_stream);
-    cudautils::device_copy_n(read_id_to_windows_section_h.data(),
-                             read_id_to_windows_section_h.size(),
-                             read_id_to_windows_section_d.data(),
-                             cuda_stream); // H2D
+    cudautils::device_copy_n_async(read_id_to_windows_section_h.data(),
+                                   read_id_to_windows_section_h.size(),
+                                   read_id_to_windows_section_d.data(),
+                                   cuda_stream); // H2D
 
     device_buffer<representation_t> window_minimizers_representation_d(total_windows, allocator, cuda_stream);
     device_buffer<char> window_minimizers_direction_d(total_windows, allocator, cuda_stream);
diff --git a/cudamapper/src/overlapper_triggered.cu b/cudamapper/src/overlapper_triggered.cu
index 331ebb936..fbc80ab29 100644
--- a/cudamapper/src/overlapper_triggered.cu
+++ b/cudamapper/src/overlapper_triggered.cu
@@ -430,7 +430,7 @@ void OverlapperTriggered::get_overlaps(std::vector<Overlap>& fused_overlaps,
 
     // memcpyD2H - move fused and filtered overlaps to host
     fused_overlaps.resize(n_filtered_overlaps);
-    cudautils::device_copy_n(d_filtered_overlaps.data(), n_filtered_overlaps, fused_overlaps.data(), _cuda_stream);
+    cudautils::device_copy_n_async(d_filtered_overlaps.data(), n_filtered_overlaps, fused_overlaps.data(), _cuda_stream);
 
     // This is not completely necessary, but if removed one has to make sure that the next step
     // uses the same stream or that sync is done in caller


Test patch:
diff --git a/cudaaligner/tests/Test_MyersAlgorithm.cu b/cudaaligner/tests/Test_MyersAlgorithm.cu
index c3829be72..ce50d5631 100644
--- a/cudaaligner/tests/Test_MyersAlgorithm.cu
+++ b/cudaaligner/tests/Test_MyersAlgorithm.cu
@@ -181,7 +181,7 @@ TEST_P(TestMyersBandedMatrixDeltas, TestCases)
     // Test if adjacent matrix entries
     // do not differ by more than delta = +/-1.
 
-    using cudautils::device_copy_n;
+    using cudautils::device_copy_n_async;
     using cudautils::set_device_value;
     using myers::word_size;
     using myers::WordType;
@@ -198,8 +198,8 @@ TEST_P(TestMyersBandedMatrixDeltas, TestCases)
     const int32_t target_size = get_size<int32_t>(t.target);
     device_buffer<char> query_d(query_size, allocator, stream.get());
     device_buffer<char> target_d(target_size, allocator, stream.get());
-    device_copy_n(t.query.c_str(), query_size, query_d.data(), stream.get());
-    device_copy_n(t.target.c_str(), target_size, target_d.data(), stream.get());
+    device_copy_n_async(t.query.c_str(), query_size, query_d.data(), stream.get());
+    device_copy_n_async(t.target.c_str(), target_size, target_d.data(), stream.get());
 
     GW_CU_CHECK_ERR(cudaStreamSynchronize(stream.get()));
 
diff --git a/cudamapper/tests/Test_CudamapperIndexCache.cu b/cudamapper/tests/Test_CudamapperIndexCache.cu
index 8f0acfb95..4e2034788 100644
--- a/cudamapper/tests/Test_CudamapperIndexCache.cu
+++ b/cudamapper/tests/Test_CudamapperIndexCache.cu
@@ -59,27 +59,27 @@ void check_if_index_is_correct(const std::shared_ptr<const Index>& index,
 
     ASSERT_EQ(get_size(index->representations()), get_size(expected_representations)) << " test_uid: " << test_uid;
     std::vector<representation_t> index_representations(index->representations().size());
-    cudautils::device_copy_n(index->representations().data(), index->representations().size(), index_representations.data(), cuda_stream); // D2H
+    cudautils::device_copy_n_async(index->representations().data(), index->representations().size(), index_representations.data(), cuda_stream); // D2H
 
     ASSERT_EQ(get_size(index->read_ids()), get_size(expected_read_ids)) << " test_uid: " << test_uid;
     std::vector<read_id_t> index_read_ids(index->read_ids().size());
-    cudautils::device_copy_n(index->read_ids().data(), index->read_ids().size(), index_read_ids.data(), cuda_stream); // D2H
+    cudautils::device_copy_n_async(index->read_ids().data(), index->read_ids().size(), index_read_ids.data(), cuda_stream); // D2H
 
     ASSERT_EQ(get_size(index->positions_in_reads()), get_size(expected_positions_in_reads)) << " test_uid: " << test_uid;
     std::vector<position_in_read_t> index_positions_in_reads(index->positions_in_reads().size());
-    cudautils::device_copy_n(index->positions_in_reads().data(), index->positions_in_reads().size(), index_positions_in_reads.data(), cuda_stream); // D2H
+    cudautils::device_copy_n_async(index->positions_in_reads().data(), index->positions_in_reads().size(), index_positions_in_reads.data(), cuda_stream); // D2H
 
     ASSERT_EQ(get_size(index->directions_of_reads()), get_size(expected_directions_of_reads)) << " test_uid: " << test_uid;
     std::vector<SketchElement::DirectionOfRepresentation> index_directions_of_reads(index->directions_of_reads().size());
-    cudautils::device_copy_n(index->directions_of_reads().data(), index->directions_of_reads().size(), index_directions_of_reads.data(), cuda_stream); // D2H
+    cudautils::device_copy_n_async(index->directions_of_reads().data(), index->directions_of_reads().size(), index_directions_of_reads.data(), cuda_stream); // D2H
 
     ASSERT_EQ(get_size(index->unique_representations()), get_size(expected_unique_representations)) << " test_uid: " << test_uid;
     std::vector<representation_t> index_unique_representations(index->unique_representations().size());
-    cudautils::device_copy_n(index->unique_representations().data(), index->unique_representations().size(), index_unique_representations.data(), cuda_stream); // D2H
+    cudautils::device_copy_n_async(index->unique_representations().data(), index->unique_representations().size(), index_unique_representations.data(), cuda_stream); // D2H
 
     ASSERT_EQ(get_size(index->first_occurrence_of_representations()), get_size(expected_first_occurrence_of_representations)) << " test_uid: " << test_uid;
     std::vector<std::uint32_t> index_first_occurrence_of_representations(index->first_occurrence_of_representations().size());
-    cudautils::device_copy_n(index->first_occurrence_of_representations().data(), index->first_occurrence_of_representations().size(), index_first_occurrence_of_representations.data(), cuda_stream); // D2H
+    cudautils::device_copy_n_async(index->first_occurrence_of_representations().data(), index->first_occurrence_of_representations().size(), index_first_occurrence_of_representations.data(), cuda_stream); // D2H
 
     read_id_t index_number_of_reads = index->number_of_reads();
 
diff --git a/cudamapper/tests/Test_CudamapperIndexGPU.cu b/cudamapper/tests/Test_CudamapperIndexGPU.cu
index 9350fcba8..b407c94ed 100644
--- a/cudamapper/tests/Test_CudamapperIndexGPU.cu
+++ b/cudamapper/tests/Test_CudamapperIndexGPU.cu
@@ -185,10 +185,10 @@ void test_find_first_occurrences_of_representations(const thrust::host_vector<re
     CudaStream cuda_stream = make_cuda_stream();
 
     device_buffer<representation_t> representations_d(representations_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(representations_h.data(),
-                             representations_h.size(),
-                             representations_d.data(),
-                             cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(representations_h.data(),
+                                   representations_h.size(),
+                                   representations_d.data(),
+                                   cuda_stream.get()); // H2D
 
     device_buffer<std::uint32_t> starting_index_of_each_representation_d(allocator, cuda_stream.get());
     device_buffer<representation_t> unique_representations_d(allocator, cuda_stream.get());
@@ -200,15 +200,15 @@ void test_find_first_occurrences_of_representations(const thrust::host_vector<re
                                               cuda_stream.get());
 
     thrust::host_vector<std::uint32_t> starting_index_of_each_representation_h(starting_index_of_each_representation_d.size());
-    cudautils::device_copy_n(starting_index_of_each_representation_d.data(),
-                             starting_index_of_each_representation_d.size(),
-                             starting_index_of_each_representation_h.data(),
-                             cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(starting_index_of_each_representation_d.data(),
+                                   starting_index_of_each_representation_d.size(),
+                                   starting_index_of_each_representation_h.data(),
+                                   cuda_stream.get()); // D2H
     thrust::host_vector<representation_t> unique_representations_h(unique_representations_d.size());
-    cudautils::device_copy_n(unique_representations_d.data(),
-                             unique_representations_d.size(),
-                             unique_representations_h.data(),
-                             cuda_stream.get()); //D2H
+    cudautils::device_copy_n_async(unique_representations_d.data(),
+                                   unique_representations_d.size(),
+                                   unique_representations_h.data(),
+                                   cuda_stream.get()); //D2H
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
 
     ASSERT_EQ(starting_index_of_each_representation_h.size(), expected_starting_index_of_each_representation_h.size());
@@ -997,17 +997,17 @@ void test_filter_out_most_common_representations(const double filtering_paramete
     CudaStream cuda_stream = make_cuda_stream();
 
     device_buffer<representation_t> representations_d(input_representations_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(input_representations_h.data(), input_representations_h.size(), representations_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(input_representations_h.data(), input_representations_h.size(), representations_d.data(), cuda_stream.get()); // H2D
     device_buffer<read_id_t> read_ids_d(input_read_ids_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(input_read_ids_h.data(), input_read_ids_h.size(), read_ids_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(input_read_ids_h.data(), input_read_ids_h.size(), read_ids_d.data(), cuda_stream.get()); // H2D
     device_buffer<position_in_read_t> positions_in_reads_d(input_positions_in_reads_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(input_positions_in_reads_h.data(), input_positions_in_reads_h.size(), positions_in_reads_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(input_positions_in_reads_h.data(), input_positions_in_reads_h.size(), positions_in_reads_d.data(), cuda_stream.get()); // H2D
     device_buffer<DirectionOfRepresentation> directions_of_representations_d(input_directions_of_representations_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(input_directions_of_representations_h.data(), input_directions_of_representations_h.size(), directions_of_representations_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(input_directions_of_representations_h.data(), input_directions_of_representations_h.size(), directions_of_representations_d.data(), cuda_stream.get()); // H2D
     device_buffer<representation_t> unique_representations_d(input_unique_representations_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(input_unique_representations_h.data(), input_unique_representations_h.size(), unique_representations_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(input_unique_representations_h.data(), input_unique_representations_h.size(), unique_representations_d.data(), cuda_stream.get()); // H2D
     device_buffer<std::uint32_t> first_occurrence_of_representations_d(input_first_occurrence_of_representations_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(input_first_occurrence_of_representations_h.data(), input_first_occurrence_of_representations_h.size(), first_occurrence_of_representations_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(input_first_occurrence_of_representations_h.data(), input_first_occurrence_of_representations_h.size(), first_occurrence_of_representations_d.data(), cuda_stream.get()); // H2D
 
     filter_out_most_common_representations(allocator,
                                            filtering_parameter,
@@ -1020,17 +1020,17 @@ void test_filter_out_most_common_representations(const double filtering_paramete
                                            cuda_stream.get());
 
     thrust::host_vector<representation_t> output_representations_h(representations_d.size());
-    cudautils::device_copy_n(representations_d.data(), representations_d.size(), output_representations_h.data(), cuda_stream.get()); //D2H
+    cudautils::device_copy_n_async(representations_d.data(), representations_d.size(), output_representations_h.data(), cuda_stream.get()); //D2H
     thrust::host_vector<read_id_t> output_read_ids_h(read_ids_d.size());
-    cudautils::device_copy_n(read_ids_d.data(), read_ids_d.size(), output_read_ids_h.data(), cuda_stream.get()); //D2H
+    cudautils::device_copy_n_async(read_ids_d.data(), read_ids_d.size(), output_read_ids_h.data(), cuda_stream.get()); //D2H
     thrust::host_vector<position_in_read_t> output_positions_in_reads_h(positions_in_reads_d.size());
-    cudautils::device_copy_n(positions_in_reads_d.data(), positions_in_reads_d.size(), output_positions_in_reads_h.data(), cuda_stream.get()); //D2H
+    cudautils::device_copy_n_async(positions_in_reads_d.data(), positions_in_reads_d.size(), output_positions_in_reads_h.data(), cuda_stream.get()); //D2H
     thrust::host_vector<DirectionOfRepresentation> output_directions_of_representations_h(directions_of_representations_d.size());
-    cudautils::device_copy_n(directions_of_representations_d.data(), directions_of_representations_d.size(), output_directions_of_representations_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(directions_of_representations_d.data(), directions_of_representations_d.size(), output_directions_of_representations_h.data(), cuda_stream.get()); // D2H
     thrust::host_vector<representation_t> output_unique_representations_h(unique_representations_d.size());
-    cudautils::device_copy_n(unique_representations_d.data(), unique_representations_d.size(), output_unique_representations_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(unique_representations_d.data(), unique_representations_d.size(), output_unique_representations_h.data(), cuda_stream.get()); // D2H
     thrust::host_vector<std::uint32_t> output_first_occurrence_of_representations_h(first_occurrence_of_representations_d.size());
-    cudautils::device_copy_n(first_occurrence_of_representations_d.data(), first_occurrence_of_representations_d.size(), output_first_occurrence_of_representations_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(first_occurrence_of_representations_d.data(), first_occurrence_of_representations_d.size(), output_first_occurrence_of_representations_h.data(), cuda_stream.get()); // D2H
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
 
     ASSERT_EQ(expected_output_representations_h.size(), output_representations_h.size());
@@ -1364,13 +1364,13 @@ void test_function(const std::string& filename,
     const device_buffer<read_id_t>& read_ids_d                                           = index.read_ids();
     const device_buffer<SketchElement::DirectionOfRepresentation>& directions_of_reads_d = index.directions_of_reads();
     thrust::host_vector<representation_t> representations_h(representations_d.size());
-    cudautils::device_copy_n(representations_d.data(), representations_d.size(), representations_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(representations_d.data(), representations_d.size(), representations_h.data(), cuda_stream.get()); // D2H
     thrust::host_vector<position_in_read_t> positions_in_reads_h(positions_in_reads_d.size());
-    cudautils::device_copy_n(positions_in_reads_d.data(), positions_in_reads_d.size(), positions_in_reads_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(positions_in_reads_d.data(), positions_in_reads_d.size(), positions_in_reads_h.data(), cuda_stream.get()); // D2H
     thrust::host_vector<read_id_t> read_ids_h(read_ids_d.size());
-    cudautils::device_copy_n(read_ids_d.data(), read_ids_d.size(), read_ids_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(read_ids_d.data(), read_ids_d.size(), read_ids_h.data(), cuda_stream.get()); // D2H
     thrust::host_vector<SketchElement::DirectionOfRepresentation> directions_of_reads_h(directions_of_reads_d.size());
-    cudautils::device_copy_n(directions_of_reads_d.data(), directions_of_reads_d.size(), directions_of_reads_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(directions_of_reads_d.data(), directions_of_reads_d.size(), directions_of_reads_h.data(), cuda_stream.get()); // D2H
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
     ASSERT_EQ(representations_h.size(), expected_representations.size());
     ASSERT_EQ(positions_in_reads_h.size(), expected_positions_in_reads.size());
@@ -1390,9 +1390,9 @@ void test_function(const std::string& filename,
     const device_buffer<representation_t>& unique_representations_d           = index.unique_representations();
     const device_buffer<std::uint32_t>& first_occurrence_of_representations_d = index.first_occurrence_of_representations();
     thrust::host_vector<representation_t> unique_representations_h(unique_representations_d.size());
-    cudautils::device_copy_n(unique_representations_d.data(), unique_representations_d.size(), unique_representations_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(unique_representations_d.data(), unique_representations_d.size(), unique_representations_h.data(), cuda_stream.get()); // D2H
     thrust::host_vector<std::uint32_t> first_occurrence_of_representations_h(first_occurrence_of_representations_d.size());
-    cudautils::device_copy_n(first_occurrence_of_representations_d.data(), first_occurrence_of_representations_d.size(), first_occurrence_of_representations_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(first_occurrence_of_representations_d.data(), first_occurrence_of_representations_d.size(), first_occurrence_of_representations_h.data(), cuda_stream.get()); // D2H
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
     ASSERT_EQ(expected_unique_representations.size() + 1, expected_first_occurrence_of_representations.size());
     ASSERT_EQ(unique_representations_h.size(), expected_unique_representations.size());
diff --git a/cudamapper/tests/Test_CudamapperMatcherGPU.cu b/cudamapper/tests/Test_CudamapperMatcherGPU.cu
index 58aa7e042..7e3ee5707 100644
--- a/cudamapper/tests/Test_CudamapperMatcherGPU.cu
+++ b/cudamapper/tests/Test_CudamapperMatcherGPU.cu
@@ -46,15 +46,15 @@ void test_find_query_target_matches(const thrust::host_vector<representation_t>&
     CudaStream cuda_stream = make_cuda_stream();
 
     device_buffer<representation_t> query_representations_d(query_representations_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(query_representations_h.data(), query_representations_h.size(), query_representations_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(query_representations_h.data(), query_representations_h.size(), query_representations_d.data(), cuda_stream.get()); // H2D
     device_buffer<representation_t> target_representations_d(target_representations_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(target_representations_h.data(), target_representations_h.size(), target_representations_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(target_representations_h.data(), target_representations_h.size(), target_representations_d.data(), cuda_stream.get()); // H2D
     device_buffer<int64_t> found_target_indices_d(query_representations_d.size(), allocator, cuda_stream.get());
 
     details::matcher_gpu::find_query_target_matches(found_target_indices_d, query_representations_d, target_representations_d, cuda_stream.get());
 
     thrust::host_vector<int64_t> found_target_indices_h(found_target_indices_d.size());
-    cudautils::device_copy_n(found_target_indices_d.data(), found_target_indices_d.size(), found_target_indices_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(found_target_indices_d.data(), found_target_indices_d.size(), found_target_indices_h.data(), cuda_stream.get()); // D2H
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
     ASSERT_EQ(found_target_indices_h.size(), expected_found_target_indices_h.size());
 
@@ -131,18 +131,18 @@ void test_compute_number_of_anchors(const thrust::host_vector<std::uint32_t>& qu
     CudaStream cuda_stream = make_cuda_stream();
 
     device_buffer<std::uint32_t> query_starting_index_of_each_representation_d(query_starting_index_of_each_representation_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(query_starting_index_of_each_representation_h.data(), query_starting_index_of_each_representation_h.size(), query_starting_index_of_each_representation_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(query_starting_index_of_each_representation_h.data(), query_starting_index_of_each_representation_h.size(), query_starting_index_of_each_representation_d.data(), cuda_stream.get()); //H2D
     device_buffer<std::uint32_t> target_starting_index_of_each_representation_d(target_starting_index_of_each_representation_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(target_starting_index_of_each_representation_h.data(), target_starting_index_of_each_representation_h.size(), target_starting_index_of_each_representation_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(target_starting_index_of_each_representation_h.data(), target_starting_index_of_each_representation_h.size(), target_starting_index_of_each_representation_d.data(), cuda_stream.get()); //H2D
     device_buffer<std::int64_t> found_target_indices_d(found_target_indices_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(found_target_indices_h.data(), found_target_indices_h.size(), found_target_indices_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(found_target_indices_h.data(), found_target_indices_h.size(), found_target_indices_d.data(), cuda_stream.get()); // H2D
     device_buffer<std::int64_t> anchor_starting_indices_d(found_target_indices_h.size(), allocator);
-    cudautils::device_copy_n(found_target_indices_h.data(), found_target_indices_h.size(), found_target_indices_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(found_target_indices_h.data(), found_target_indices_h.size(), found_target_indices_d.data(), cuda_stream.get()); // H2D
 
     details::matcher_gpu::compute_anchor_starting_indices(anchor_starting_indices_d, query_starting_index_of_each_representation_d, found_target_indices_d, target_starting_index_of_each_representation_d, cuda_stream.get());
 
     thrust::host_vector<std::int64_t> anchor_starting_indices_h(anchor_starting_indices_d.size());
-    cudautils::device_copy_n(anchor_starting_indices_d.data(), anchor_starting_indices_d.size(), anchor_starting_indices_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(anchor_starting_indices_d.data(), anchor_starting_indices_d.size(), anchor_starting_indices_h.data(), cuda_stream.get()); // D2H
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
     for (int32_t i = 0; i < get_size(found_target_indices_h); ++i)
     {
@@ -240,21 +240,21 @@ void test_generate_anchors(
     CudaStream cuda_stream = make_cuda_stream();
 
     device_buffer<std::int64_t> anchor_starting_indices_d(anchor_starting_indices_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(anchor_starting_indices_h.data(), anchor_starting_indices_h.size(), anchor_starting_indices_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(anchor_starting_indices_h.data(), anchor_starting_indices_h.size(), anchor_starting_indices_d.data(), cuda_stream.get()); // H2D
     device_buffer<std::uint32_t> query_starting_index_of_each_representation_d(query_starting_index_of_each_representation_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(query_starting_index_of_each_representation_h.data(), query_starting_index_of_each_representation_h.size(), query_starting_index_of_each_representation_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(query_starting_index_of_each_representation_h.data(), query_starting_index_of_each_representation_h.size(), query_starting_index_of_each_representation_d.data(), cuda_stream.get()); // H2D
     device_buffer<std::int64_t> found_target_indices_d(found_target_indices_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(found_target_indices_h.data(), found_target_indices_h.size(), found_target_indices_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(found_target_indices_h.data(), found_target_indices_h.size(), found_target_indices_d.data(), cuda_stream.get()); //H2D
     device_buffer<std::uint32_t> target_starting_index_of_each_representation_d(target_starting_index_of_each_representation_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(target_starting_index_of_each_representation_h.data(), target_starting_index_of_each_representation_h.size(), target_starting_index_of_each_representation_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(target_starting_index_of_each_representation_h.data(), target_starting_index_of_each_representation_h.size(), target_starting_index_of_each_representation_d.data(), cuda_stream.get()); // H2D
     device_buffer<read_id_t> query_read_ids_d(query_read_ids_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(query_read_ids_h.data(), query_read_ids_h.size(), query_read_ids_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(query_read_ids_h.data(), query_read_ids_h.size(), query_read_ids_d.data(), cuda_stream.get()); // H2D
     device_buffer<position_in_read_t> query_positions_in_read_d(query_positions_in_read_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(query_positions_in_read_h.data(), query_positions_in_read_h.size(), query_positions_in_read_d.data(), cuda_stream.get()); // H2D
+    cudautils::device_copy_n_async(query_positions_in_read_h.data(), query_positions_in_read_h.size(), query_positions_in_read_d.data(), cuda_stream.get()); // H2D
     device_buffer<read_id_t> target_read_ids_d(target_read_ids_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(target_read_ids_h.data(), target_read_ids_h.size(), target_read_ids_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(target_read_ids_h.data(), target_read_ids_h.size(), target_read_ids_d.data(), cuda_stream.get()); //H2D
     device_buffer<position_in_read_t> target_positions_in_read_d(target_positions_in_read_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(target_positions_in_read_h.data(), target_positions_in_read_h.size(), target_positions_in_read_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(target_positions_in_read_h.data(), target_positions_in_read_h.size(), target_positions_in_read_d.data(), cuda_stream.get()); //H2D
 
     device_buffer<Anchor> anchors_d(anchor_starting_indices_h.back(), allocator, cuda_stream.get());
 
@@ -281,7 +281,7 @@ void test_generate_anchors(
                                                       target_index);
 
     thrust::host_vector<Anchor> anchors_h(anchors_d.size());
-    cudautils::device_copy_n(anchors_d.data(), anchors_d.size(), anchors_h.data(), cuda_stream.get()); // D2H
+    cudautils::device_copy_n_async(anchors_d.data(), anchors_d.size(), anchors_h.data(), cuda_stream.get()); // D2H
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
     ASSERT_EQ(anchors_h.size(), expected_anchors_h.size());
 
diff --git a/cudamapper/tests/Test_CudamapperMinimizer.cpp b/cudamapper/tests/Test_CudamapperMinimizer.cpp
index 3323edfc5..f252aa608 100644
--- a/cudamapper/tests/Test_CudamapperMinimizer.cpp
+++ b/cudamapper/tests/Test_CudamapperMinimizer.cpp
@@ -43,16 +43,16 @@ void test_function(const std::uint64_t number_of_reads_to_add,
     CudaStream cuda_stream = make_cuda_stream();
 
     device_buffer<char> merged_basepairs_d(merged_basepairs_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(merged_basepairs_h.data(),
-                             merged_basepairs_h.size(),
-                             merged_basepairs_d.data(),
-                             cuda_stream.get());
+    cudautils::device_copy_n_async(merged_basepairs_h.data(),
+                                   merged_basepairs_h.size(),
+                                   merged_basepairs_d.data(),
+                                   cuda_stream.get());
 
     device_buffer<ArrayBlock> read_id_to_basepairs_section_d(read_id_to_basepairs_section_h.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(read_id_to_basepairs_section_h.data(),
-                             read_id_to_basepairs_section_h.size(),
-                             read_id_to_basepairs_section_d.data(),
-                             cuda_stream.get());
+    cudautils::device_copy_n_async(read_id_to_basepairs_section_h.data(),
+                                   read_id_to_basepairs_section_h.size(),
+                                   read_id_to_basepairs_section_d.data(),
+                                   cuda_stream.get());
 
     auto sketch_elements = Minimizer::generate_sketch_elements(allocator,
                                                                number_of_reads_to_add,
@@ -67,15 +67,15 @@ void test_function(const std::uint64_t number_of_reads_to_add,
 
     device_buffer<representation_t> representations_d = std::move(sketch_elements.representations_d);
     std::vector<representation_t> representations_h(representations_d.size());
-    cudautils::device_copy_n(representations_d.data(),
-                             representations_d.size(),
-                             representations_h.data(),
-                             cuda_stream.get());
+    cudautils::device_copy_n_async(representations_d.data(),
+                                   representations_d.size(),
+                                   representations_h.data(),
+                                   cuda_stream.get());
     device_buffer<Minimizer::ReadidPositionDirection> rest_d = std::move(sketch_elements.rest_d);
     std::vector<Minimizer::ReadidPositionDirection> rest_h(rest_d.size());
-    cudautils::device_copy_n(rest_d.data(),
-                             rest_d.size(),
-                             rest_h.data(), cuda_stream.get());
+    cudautils::device_copy_n_async(rest_d.data(),
+                                   rest_d.size(),
+                                   rest_h.data(), cuda_stream.get());
     GW_CU_CHECK_ERR(cudaStreamSynchronize(cuda_stream.get()));
 
     ASSERT_EQ(expected_representations_h.size(), expected_rest_h.size());
diff --git a/cudamapper/tests/Test_CudamapperOverlapperTriggered.cu b/cudamapper/tests/Test_CudamapperOverlapperTriggered.cu
index cab478b2e..bac8d92ec 100644
--- a/cudamapper/tests/Test_CudamapperOverlapperTriggered.cu
+++ b/cudamapper/tests/Test_CudamapperOverlapperTriggered.cu
@@ -46,7 +46,7 @@ TEST(TestCudamapperOverlapperTriggerred, OneAchorNoOverlaps)
     anchors.push_back(anchor1);
 
     device_buffer<Anchor> anchors_d(anchors.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
 
     std::vector<Overlap> overlaps;
     overlapper.get_overlaps(overlaps, anchors_d, 0, 0);
@@ -93,7 +93,7 @@ TEST(TestCudamapperOverlapperTriggerred, FourAnchorsOneOverlap)
     anchors.push_back(anchor4);
 
     device_buffer<Anchor> anchors_d(anchors.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
 
     std::vector<Overlap> overlaps;
     overlapper.get_overlaps(overlaps, anchors_d, false, 0, 0, 1000);
@@ -146,7 +146,7 @@ TEST(TestCudamapperOverlapperTriggerred, FourAnchorsNoOverlap)
     anchors.push_back(anchor4);
 
     device_buffer<Anchor> anchors_d(anchors.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
 
     std::vector<Overlap> overlaps;
     overlapper.get_overlaps(overlaps, anchors_d, 0, 0, 1000);
@@ -193,7 +193,7 @@ TEST(TestCudamapperOverlapperTriggerred, FourColinearAnchorsOneOverlap)
     anchors.push_back(anchor4);
 
     device_buffer<Anchor> anchors_d(anchors.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
 
     std::vector<Overlap> overlaps;
     overlapper.get_overlaps(overlaps, anchors_d, 0, 0);
@@ -240,7 +240,7 @@ TEST(TestCudamapperOverlapperTriggerred, FourAnchorsLastNotInOverlap)
     anchors.push_back(anchor4);
 
     device_buffer<Anchor> anchors_d(anchors.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
 
     std::vector<Overlap> overlaps;
     overlapper.get_overlaps(overlaps, anchors_d, false, 0, 0, 1000);
@@ -293,7 +293,7 @@ TEST(TestCudamapperOverlapperTriggerred, ReverseStrand)
     anchors.push_back(anchor4);
 
     device_buffer<Anchor> anchors_d(anchors.size(), allocator, cuda_stream.get());
-    cudautils::device_copy_n(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
+    cudautils::device_copy_n_async(anchors.data(), anchors.size(), anchors_d.data(), cuda_stream.get()); //H2D
 
     std::vector<Overlap> overlaps;
     overlapper.get_overlaps(overlaps, anchors_d, false, 0, 0, 1000);