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

Issue:
[cudaaligner] use device_copy_n
cudaaligner uses many `cudaMemcpy`, many (if not all) can be replaced with `device_copy_n`.


Discussion / hints:

Fix patch:
diff --git a/cudaaligner/src/aligner_global.cpp b/cudaaligner/src/aligner_global.cpp
index 59892929a..b6a797b1c 100644
--- a/cudaaligner/src/aligner_global.cpp
+++ b/cudaaligner/src/aligner_global.cpp
@@ -149,16 +149,8 @@ StatusType AlignerGlobal::align_all()
     scoped_device_switch dev(device_id_);
     const int32_t max_alignment_length = std::max(max_query_length_, max_target_length_);
     const int32_t max_result_length    = calc_max_result_length(max_query_length_, max_target_length_);
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequence_lengths_d_.data(),
-                                    sequence_lengths_h_.data(),
-                                    2 * sizeof(int32_t) * num_alignments,
-                                    cudaMemcpyHostToDevice,
-                                    stream_));
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequences_d_.data(),
-                                    sequences_h_.data(),
-                                    2 * sizeof(char) * max_alignment_length * num_alignments,
-                                    cudaMemcpyHostToDevice,
-                                    stream_));
+    cudautils::device_copy_n_async(sequence_lengths_h_.data(), 2 * num_alignments, sequence_lengths_d_.data(), stream_);
+    cudautils::device_copy_n_async(sequences_h_.data(), 2 * max_alignment_length * num_alignments, sequences_d_.data(), stream_);
 
     // Run kernel
     run_alignment(results_d_.data(), result_lengths_d_.data(),
@@ -167,16 +159,8 @@ StatusType AlignerGlobal::align_all()
                   num_alignments,
                   stream_);
 
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(results_h_.data(),
-                                    results_d_.data(),
-                                    sizeof(int8_t) * max_result_length * num_alignments,
-                                    cudaMemcpyDeviceToHost,
-                                    stream_));
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(result_lengths_h_.data(),
-                                    result_lengths_d_.data(),
-                                    sizeof(int32_t) * num_alignments,
-                                    cudaMemcpyDeviceToHost,
-                                    stream_));
+    cudautils::device_copy_n_async(results_d_.data(), max_result_length * num_alignments, results_h_.data(), stream_);
+    cudautils::device_copy_n_async(result_lengths_d_.data(), num_alignments, result_lengths_h_.data(), stream_);
     return StatusType::success;
 }
 
diff --git a/cudaaligner/src/myers_gpu.cu b/cudaaligner/src/myers_gpu.cu
index aebb47d87..a7a9f1dad 100644
--- a/cudaaligner/src/myers_gpu.cu
+++ b/cudaaligner/src/myers_gpu.cu
@@ -904,9 +904,9 @@ int32_t myers_compute_edit_distance(std::string const& target, std::string const
     batched_device_matrices<myers::WordType> query_patterns(1, n_words * 4, allocator, stream.get());
 
     std::array<int32_t, 2> lengths = {static_cast<int32_t>(get_size(query)), static_cast<int32_t>(get_size(target))};
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequences_d.data(), query.data(), sizeof(char) * get_size(query), cudaMemcpyHostToDevice, stream.get()));
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequences_d.data() + max_sequence_length, target.data(), sizeof(char) * get_size(target), cudaMemcpyHostToDevice, stream.get()));
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequence_lengths_d.data(), lengths.data(), sizeof(int32_t) * 2, cudaMemcpyHostToDevice, stream.get()));
+    cudautils::device_copy_n_async(query.data(), get_size(query), sequences_d.data(), stream.get());
+    cudautils::device_copy_n_async(target.data(), get_size(target), sequences_d.data() + max_sequence_length, stream.get());
+    cudautils::device_copy_n_async(lengths.data(), 2, sequence_lengths_d.data(), stream.get());
 
     myers::myers_compute_score_matrix_kernel<<<1, warp_size, 0, stream.get()>>>(pv.get_device_interface(), mv.get_device_interface(), score.get_device_interface(), query_patterns.get_device_interface(), sequences_d.data(), sequence_lengths_d.data(), max_sequence_length, 1);
     GW_CU_CHECK_ERR(cudaPeekAtLastError());
@@ -949,9 +949,9 @@ matrix<int32_t> myers_get_full_score_matrix(std::string const& target, std::stri
     batched_device_matrices<int32_t> fullscore(1, (get_size(query) + 1) * (get_size(target) + 1), allocator, stream.get());
 
     std::array<int32_t, 2> lengths = {static_cast<int32_t>(get_size(query)), static_cast<int32_t>(get_size(target))};
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequences_d.data(), query.data(), sizeof(char) * get_size(query), cudaMemcpyHostToDevice, stream.get()));
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequences_d.data() + max_sequence_length, target.data(), sizeof(char) * get_size(target), cudaMemcpyHostToDevice, stream.get()));
-    GW_CU_CHECK_ERR(cudaMemcpyAsync(sequence_lengths_d.data(), lengths.data(), sizeof(int32_t) * 2, cudaMemcpyHostToDevice, stream.get()));
+    cudautils::device_copy_n_async(query.data(), get_size(query), sequences_d.data(), stream.get());
+    cudautils::device_copy_n_async(target.data(), get_size(target), sequences_d.data() + max_sequence_length, stream.get());
+    cudautils::device_copy_n_async(lengths.data(), 2, sequence_lengths_d.data(), stream.get());
 
     myers::myers_compute_score_matrix_kernel<<<1, warp_size, 0, stream.get()>>>(pv.get_device_interface(), mv.get_device_interface(), score.get_device_interface(), query_patterns.get_device_interface(), sequences_d.data(), sequence_lengths_d.data(), max_sequence_length, 1);
     GW_CU_CHECK_ERR(cudaPeekAtLastError());


Test patch:
diff --git a/cudaaligner/tests/Test_HirschbergMyers.cu b/cudaaligner/tests/Test_HirschbergMyers.cu
index dff4db82e..a997e8be7 100644
--- a/cudaaligner/tests/Test_HirschbergMyers.cu
+++ b/cudaaligner/tests/Test_HirschbergMyers.cu
@@ -60,7 +60,7 @@ matrix<WordType> compute_myers_preprocess_matrix(std::string query_host)
     const int32_t n_words            = ceiling_divide<int32_t>(query_host.size(), word_size);
 
     device_buffer<char> query(query_host.size(), allocator);
-    GW_CU_CHECK_ERR(cudaMemcpy(query.data(), query_host.data(), sizeof(char) * query.size(), cudaMemcpyHostToDevice));
+    cudautils::device_copy_n(query_host.data(), query.size(), query.data());
 
     batched_device_matrices<WordType> query_pattern(1, 8 * n_words, allocator, stream);
     myers_preprocess_kernel<<<1, 32>>>(query_pattern.get_device_interface(), query.data(), query.size());
@@ -75,7 +75,7 @@ std::vector<WordType> myers_get_query_pattern_test(std::string query_host, int32
     cudaStream_t stream              = nullptr;
     const int32_t n_words            = ceiling_divide<int32_t>(query_host.size(), word_size);
     device_buffer<char> query(query_host.size(), allocator);
-    GW_CU_CHECK_ERR(cudaMemcpy(query.data(), query_host.data(), sizeof(char) * query.size(), cudaMemcpyHostToDevice));
+    cudautils::device_copy_n(query_host.data(), query.size(), query.data());
     batched_device_matrices<WordType> query_pattern(1, 8 * n_words, allocator, stream);
     myers_preprocess_kernel<<<1, 32>>>(query_pattern.get_device_interface(), query.data(), query.size());
     GW_CU_CHECK_ERR(cudaPeekAtLastError());
@@ -85,7 +85,7 @@ std::vector<WordType> myers_get_query_pattern_test(std::string query_host, int32
     GW_CU_CHECK_ERR(cudaPeekAtLastError());
 
     std::vector<WordType> result_host(result.size());
-    cudaMemcpy(result_host.data(), result.data(), sizeof(WordType) * result.size(), cudaMemcpyDeviceToHost);
+    cudautils::device_copy_n(result.data(), result.size(), result_host.data());
     return result_host;
 }
 
diff --git a/cudaaligner/tests/Test_NeedlemanWunschImplementation.cpp b/cudaaligner/tests/Test_NeedlemanWunschImplementation.cpp
index ca5b76302..1e84edc5e 100644
--- a/cudaaligner/tests/Test_NeedlemanWunschImplementation.cpp
+++ b/cudaaligner/tests/Test_NeedlemanWunschImplementation.cpp
@@ -172,12 +172,12 @@ matrix<int> ukkonen_gpu_build_score_matrix(const std::string& target, const std:
     std::vector<int32_t> path_length_h(1);
 
     device_buffer<char> sequences_d(2 * max_alignment_length, allocator);
-    GW_CU_CHECK_ERR(cudaMemcpy(sequences_d.data(), query.c_str(), sizeof(char) * query_length, cudaMemcpyHostToDevice));
-    GW_CU_CHECK_ERR(cudaMemcpy(sequences_d.data() + max_alignment_length, target.c_str(), sizeof(char) * target_length, cudaMemcpyHostToDevice));
+    cudautils::device_copy_n(query.c_str(), query_length, sequences_d.data());
+    cudautils::device_copy_n(target.c_str(), target_length, sequences_d.data() + max_alignment_length);
 
     device_buffer<int32_t> sequence_lengths_d(2, allocator);
-    GW_CU_CHECK_ERR(cudaMemcpy(sequence_lengths_d.data(), &query_length, sizeof(int32_t) * 1, cudaMemcpyHostToDevice));
-    GW_CU_CHECK_ERR(cudaMemcpy(sequence_lengths_d.data() + 1, &target_length, sizeof(int32_t) * 1, cudaMemcpyHostToDevice));
+    cudautils::set_device_value(sequence_lengths_d.data(), query_length);
+    cudautils::set_device_value(sequence_lengths_d.data() + 1, target_length);
 
     ukkonen_compute_score_matrix_gpu(*score_matrices.get(),
                                      sequences_d.data(), sequence_lengths_d.data(),
@@ -253,12 +253,12 @@ std::vector<int8_t> run_ukkonen_gpu(const std::string& target, const std::string
     std::vector<int32_t> path_length_h(1);
 
     device_buffer<char> sequences_d(2 * max_alignment_length, allocator);
-    GW_CU_CHECK_ERR(cudaMemcpy(sequences_d.data(), query.c_str(), sizeof(char) * query_length, cudaMemcpyHostToDevice));
-    GW_CU_CHECK_ERR(cudaMemcpy(sequences_d.data() + max_alignment_length, target.c_str(), sizeof(char) * target_length, cudaMemcpyHostToDevice));
+    cudautils::device_copy_n(query.c_str(), query_length, sequences_d.data());
+    cudautils::device_copy_n(target.c_str(), target_length, sequences_d.data() + max_alignment_length);
 
     device_buffer<int32_t> sequence_lengths_d(2, allocator);
-    GW_CU_CHECK_ERR(cudaMemcpy(sequence_lengths_d.data(), &query_length, sizeof(int32_t) * 1, cudaMemcpyHostToDevice));
-    GW_CU_CHECK_ERR(cudaMemcpy(sequence_lengths_d.data() + 1, &target_length, sizeof(int32_t) * 1, cudaMemcpyHostToDevice));
+    cudautils::set_device_value(sequence_lengths_d.data(), query_length);
+    cudautils::set_device_value(sequence_lengths_d.data() + 1, target_length);
 
     // Run kernel
     ukkonen_gpu(path_d.data(), path_length_d.data(), max_path_length,
@@ -269,8 +269,8 @@ std::vector<int8_t> run_ukkonen_gpu(const std::string& target, const std::string
                 nullptr);
 
     // Get results
-    GW_CU_CHECK_ERR(cudaMemcpy(path_h.data(), path_d.data(), sizeof(int8_t) * max_path_length, cudaMemcpyDeviceToHost));
-    GW_CU_CHECK_ERR(cudaMemcpy(path_length_h.data(), path_length_d.data(), sizeof(int32_t) * 1, cudaMemcpyDeviceToHost));
+    cudautils::device_copy_n(path_d.data(), max_path_length, path_h.data());
+    cudautils::device_copy_n(path_length_d.data(), 1, path_length_h.data());
 
     std::vector<int8_t> bt;
     for (int32_t l = 0; l < path_length_h[0]; l++)