From 0ab2f735cda2fdd3060cbf2ea8af23b2d896a2ba Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Mon, 10 Aug 2026 09:11:52 -0500 Subject: [PATCH 1/3] Require cudaMemcpyDefault policy --- .pre-commit-config.yaml | 6 ++++++ cpp/include/cudf_test/tdigest_utilities.hpp | 8 +++----- cpp/src/io/utilities/data_sink.cpp | 2 +- cpp/src/strings/utilities.cu | 8 ++++---- cpp/tests/io/cudftable_test.cpp | 2 +- cpp/tests/io/io_test_utils.hpp | 2 +- cpp/tests/io/json/json_quote_normalization_test.cpp | 4 ++-- java/src/main/native/src/jni_writer_data_sink.hpp | 6 +++--- java/src/main/native/src/multi_host_buffer_source.cpp | 4 ++-- python/pylibcudf/pylibcudf/contiguous_split.pyx | 2 +- .../code/threaded_handoff.cu | 2 +- 11 files changed, 25 insertions(+), 21 deletions(-) diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 7e6924f1f271..fd901f588d7c 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -119,6 +119,12 @@ repos: language: pygrep types_or: [c, c++, cuda] files: '^cpp/(src|include)/' + - id: use-cuda-memcpy-default + name: use-cuda-memcpy-default + description: 'Enforce that cudaMemcpyDefault is used instead of explicit host/device cudaMemcpyKind policies' + entry: '\bcudaMemcpy(?:HostToHost|HostToDevice|DeviceToHost|DeviceToDevice)\b' + language: pygrep + types_or: [c, c++, cuda, cython] - id: use-cudf-memcpy-async name: use-cudf-memcpy-async description: 'Enforce that cudf::detail::memcpy_async or memcpy_batch_async is used instead of cudaMemcpyAsync (see developer guide)' diff --git a/cpp/include/cudf_test/tdigest_utilities.hpp b/cpp/include/cudf_test/tdigest_utilities.hpp index dc252d14b08a..e6013c9abf9c 100644 --- a/cpp/include/cudf_test/tdigest_utilities.hpp +++ b/cpp/include/cudf_test/tdigest_utilities.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -128,10 +128,8 @@ void tdigest_minmax_compare(cudf::tdigest::tdigest_column_view const& tdv, auto expected_max = static_cast(max_scalar->value()); double tdv_min, tdv_max; - EXPECT_EQ(cudaMemcpy(&tdv_min, tdv.min_begin(), sizeof(double), cudaMemcpyDeviceToHost), - cudaSuccess); - EXPECT_EQ(cudaMemcpy(&tdv_max, tdv.max_begin(), sizeof(double), cudaMemcpyDeviceToHost), - cudaSuccess); + EXPECT_EQ(cudaMemcpy(&tdv_min, tdv.min_begin(), sizeof(double), cudaMemcpyDefault), cudaSuccess); + EXPECT_EQ(cudaMemcpy(&tdv_max, tdv.max_begin(), sizeof(double), cudaMemcpyDefault), cudaSuccess); EXPECT_EQ(tdv_min, expected_min); EXPECT_EQ(tdv_max, expected_max); diff --git a/cpp/src/io/utilities/data_sink.cpp b/cpp/src/io/utilities/data_sink.cpp index d2f89c454c79..f77ca5580ea7 100644 --- a/cpp/src/io/utilities/data_sink.cpp +++ b/cpp/src/io/utilities/data_sink.cpp @@ -120,7 +120,7 @@ class host_buffer_sink : public data_sink { // in-flight when using cudaMemcpySrcAccessOrderStream. Need to ensure // stream ordering or pre-reserve buffer to avoid reallocation. CUDF_CUDA_TRY(cudaMemcpyAsync( - buffer_->data() + current_size, gpu_data, size, cudaMemcpyDeviceToHost, stream.value())); + buffer_->data() + current_size, gpu_data, size, cudaMemcpyDefault, stream.value())); return std::async(std::launch::deferred, [stream]() -> void { stream.synchronize(); }); } diff --git a/cpp/src/strings/utilities.cu b/cpp/src/strings/utilities.cu index 8afb7eea7661..8d256b032944 100644 --- a/cpp/src/strings/utilities.cu +++ b/cpp/src/strings/utilities.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -106,7 +106,7 @@ character_flags_table_type const* get_character_flags_table(rmm::cuda_stream_vie g_character_codepoint_flags, sizeof(g_character_codepoint_flags), 0, - cudaMemcpyHostToDevice, + cudaMemcpyDefault, stream.value())); CUDF_CUDA_TRY(cudaGetSymbolAddress((void**)&table, character_codepoint_flags)); return table; @@ -124,7 +124,7 @@ character_cases_table_type const* get_character_cases_table(rmm::cuda_stream_vie g_character_cases_table, sizeof(g_character_cases_table), 0, - cudaMemcpyHostToDevice, + cudaMemcpyDefault, stream.value())); CUDF_CUDA_TRY(cudaGetSymbolAddress((void**)&table, character_cases_table)); return table; @@ -142,7 +142,7 @@ special_case_mapping const* get_special_case_mapping_table(rmm::cuda_stream_view g_special_case_mappings, sizeof(g_special_case_mappings), 0, - cudaMemcpyHostToDevice, + cudaMemcpyDefault, stream.value())); CUDF_CUDA_TRY(cudaGetSymbolAddress((void**)&table, character_special_case_mappings)); return table; diff --git a/cpp/tests/io/cudftable_test.cpp b/cpp/tests/io/cudftable_test.cpp index 63859cb99d30..8663827dfd75 100644 --- a/cpp/tests/io/cudftable_test.cpp +++ b/cpp/tests/io/cudftable_test.cpp @@ -515,7 +515,7 @@ TEST_F(CudftableTest, DeviceBufferSource) rmm::device_buffer device_buffer(buffer.size(), cudf::get_default_stream()); auto const stream = cudf::get_default_stream(); CUDF_CUDA_TRY(cudaMemcpyAsync( - device_buffer.data(), buffer.data(), buffer.size(), cudaMemcpyHostToDevice, stream.value())); + device_buffer.data(), buffer.data(), buffer.size(), cudaMemcpyDefault, stream.value())); // Ensure the data is copied to the device before the host read, because the host read does not // take the stream stream.synchronize(); diff --git a/cpp/tests/io/io_test_utils.hpp b/cpp/tests/io/io_test_utils.hpp index 9fb6c64bb108..c91317eed6b4 100644 --- a/cpp/tests/io/io_test_utils.hpp +++ b/cpp/tests/io/io_test_utils.hpp @@ -60,7 +60,7 @@ class ThrowingDeviceReadDatasource : public cudf::io::datasource { size = std::min(size, data_.size() - offset); rmm::device_buffer out_data(size, stream); cudaMemcpyAsync( - out_data.data(), data_.data() + offset, size, cudaMemcpyHostToDevice, stream.value()); + out_data.data(), data_.data() + offset, size, cudaMemcpyDefault, stream.value()); cudaStreamSynchronize(stream.value()); return cudf::io::datasource::buffer::create(std::move(out_data)); } diff --git a/cpp/tests/io/json/json_quote_normalization_test.cpp b/cpp/tests/io/json/json_quote_normalization_test.cpp index d3c8c944e8fe..760fa3a5e255 100644 --- a/cpp/tests/io/json/json_quote_normalization_test.cpp +++ b/cpp/tests/io/json/json_quote_normalization_test.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -37,7 +37,7 @@ void run_test(std::string const& host_input, CUDF_CUDA_TRY(cudaMemcpyAsync(preprocessed_host_output.data(), device_data.data(), preprocessed_host_output.size(), - cudaMemcpyDeviceToHost, + cudaMemcpyDefault, stream_view.value())) stream_view.synchronize(); CUDF_TEST_EXPECT_VECTOR_EQUAL( diff --git a/java/src/main/native/src/jni_writer_data_sink.hpp b/java/src/main/native/src/jni_writer_data_sink.hpp index d102c51eff6f..d5e9e0bb7ab4 100644 --- a/java/src/main/native/src/jni_writer_data_sink.hpp +++ b/java/src/main/native/src/jni_writer_data_sink.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -89,8 +89,8 @@ class jni_writer_data_sink final : public cudf::io::data_sink { left_to_copy < buffer_amount_available ? left_to_copy : buffer_amount_available; char* copy_to = current_buffer_data + current_buffer_written; - CUDF_CUDA_TRY(cudaMemcpyAsync( - copy_to, copy_from, amount_to_copy, cudaMemcpyDeviceToHost, stream.value())); + CUDF_CUDA_TRY( + cudaMemcpyAsync(copy_to, copy_from, amount_to_copy, cudaMemcpyDefault, stream.value())); copy_from = copy_from + amount_to_copy; current_buffer_written += amount_to_copy; diff --git a/java/src/main/native/src/multi_host_buffer_source.cpp b/java/src/main/native/src/multi_host_buffer_source.cpp index abc4cc806c8c..6e18f3ea1de9 100644 --- a/java/src/main/native/src/multi_host_buffer_source.cpp +++ b/java/src/main/native/src/multi_host_buffer_source.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -114,7 +114,7 @@ size_t multi_host_buffer_source::device_read(size_t offset, auto buffer_offset = offset - offsets_[buffer_index]; auto src = addrs_[buffer_index] + buffer_offset; auto copy_size = std::min(buffer_left, bytes_left); - CUDF_CUDA_TRY(cudaMemcpyAsync(dst, src, copy_size, cudaMemcpyHostToDevice, stream.value())); + CUDF_CUDA_TRY(cudaMemcpyAsync(dst, src, copy_size, cudaMemcpyDefault, stream.value())); offset += copy_size; dst += copy_size; bytes_left -= copy_size; diff --git a/python/pylibcudf/pylibcudf/contiguous_split.pyx b/python/pylibcudf/pylibcudf/contiguous_split.pyx index 1f5cb5c52f89..b2802835fe21 100644 --- a/python/pylibcudf/pylibcudf/contiguous_split.pyx +++ b/python/pylibcudf/pylibcudf/contiguous_split.pyx @@ -303,7 +303,7 @@ cdef class ChunkedPack: dereference(h_buf).data() + offset, d_span.data(), size, - cudaMemcpyKind.cudaMemcpyDeviceToHost, + cudaMemcpyKind.cudaMemcpyDefault, stream, ) offset += size diff --git a/skills/accelerated-computing-cudf/evals/files/cudf-native-stream-handoff-boundary/code/threaded_handoff.cu b/skills/accelerated-computing-cudf/evals/files/cudf-native-stream-handoff-boundary/code/threaded_handoff.cu index 6e0284c72fc6..270e0e3b5d52 100644 --- a/skills/accelerated-computing-cudf/evals/files/cudf-native-stream-handoff-boundary/code/threaded_handoff.cu +++ b/skills/accelerated-computing-cudf/evals/files/cudf-native-stream-handoff-boundary/code/threaded_handoff.cu @@ -96,7 +96,7 @@ std::uint64_t consume_on_stream(const std::shared_ptr& table, check_cuda(cudaMemcpyAsync(&h_sum, d_sum, sizeof(std::uint64_t), - cudaMemcpyDeviceToHost, + cudaMemcpyDefault, consumer_stream), "copy checksum"); check_cuda(cudaStreamSynchronize(consumer_stream), "sync consumer stream"); From 121e69dbde58f2fdaf64afea7aed02ef717ef0c4 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Mon, 10 Aug 2026 10:03:42 -0500 Subject: [PATCH 2/3] Update developer guide --- cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md index 8c7f97360123..0d5e5d6dcca5 100644 --- a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md +++ b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md @@ -882,6 +882,13 @@ temporary host staging buffers to avoid the sync: The same stream-safety requirements apply to `memcpy_async` and `memcpy_batch_async`. +If CUDA memory copy APIs must be called directly, always use `cudaMemcpyDefault` instead of an +explicit host/device copy policy. Copy correctness depends on whether the source and destination +pointers are accessible from the host or device, not where the memory is resident. For example, +pinned host memory may be device-accessible despite residing on the host. `cudaMemcpyDefault` allows +CUDA to infer the valid copy direction from the pointers rather than rejecting such copies based on +an explicit policy. + ## Default Parameters While public libcudf APIs are free to include default function parameters, detail functions should From 4c989bc96edf721b95b6f1d46df7ad7e32f4c0ee Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Mon, 10 Aug 2026 10:04:16 -0500 Subject: [PATCH 3/3] Use CUDF_CUDA_TRY --- cpp/tests/io/io_test_utils.hpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/cpp/tests/io/io_test_utils.hpp b/cpp/tests/io/io_test_utils.hpp index c91317eed6b4..30946b43603f 100644 --- a/cpp/tests/io/io_test_utils.hpp +++ b/cpp/tests/io/io_test_utils.hpp @@ -8,6 +8,7 @@ #include #include #include +#include #include @@ -59,9 +60,9 @@ class ThrowingDeviceReadDatasource : public cudf::io::datasource { // For testing, just copy the data from the host buffer into a new buffer size = std::min(size, data_.size() - offset); rmm::device_buffer out_data(size, stream); - cudaMemcpyAsync( - out_data.data(), data_.data() + offset, size, cudaMemcpyDefault, stream.value()); - cudaStreamSynchronize(stream.value()); + CUDF_CUDA_TRY(cudaMemcpyAsync( + out_data.data(), data_.data() + offset, size, cudaMemcpyDefault, stream.value())); + stream.synchronize(); return cudf::io::datasource::buffer::create(std::move(out_data)); }