Require cudaMemcpyDefault policy

.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)'

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

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<double>(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);

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(); });
   }
 

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;

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();

cpp/tests/io/io_test_utils.hpp

@@ -8,6 +8,7 @@
 #include <cudf/io/data_sink.hpp>
 #include <cudf/io/datasource.hpp>
 #include <cudf/logger.hpp>
+#include <cudf/utilities/error.hpp>
 
 #include <rapids_logger/logger.hpp>
 
@@ -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, cudaMemcpyHostToDevice, 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));
   }
 

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(

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;

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;

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

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<NativeGpuTable>& 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");