diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml
index 500f5a71636b..5abeabc39691 100644
--- a/.github/workflows/pr.yaml
+++ b/.github/workflows/pr.yaml
@@ -46,7 +46,7 @@ jobs:
       - narwhals-tests
       - telemetry-setup
       - third-party-integration-tests-cudf-pandas
-    uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@test-gb300
     permissions:
       contents: read
     if: always()
@@ -89,7 +89,7 @@ jobs:
       packages: read
       pull-requests: read
     needs: telemetry-setup
-    uses: rapidsai/shared-workflows/.github/workflows/changed-files.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/changed-files.yaml@test-gb300
     with:
       files_yaml: |
         build_docs:
@@ -367,7 +367,7 @@ jobs:
     permissions:
       contents: read
     needs: telemetry-setup
-    uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@test-gb300
     with:
       enable_check_generated_files: false
       ignored_pr_jobs: "telemetry-summarize spark-rapids-jni cuml-compat-tests"
@@ -380,7 +380,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@test-gb300
     with:
       build_type: pull-request
       build-datetime: ${{ needs.build-details.outputs.build-datetime }}
@@ -395,7 +395,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     with:
       build_type: pull-request
       node_type: cpu16
@@ -411,7 +411,7 @@ jobs:
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
     needs: checks
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     with:
       build_type: pull-request
       script: "ci/cpp_linters.sh"
@@ -423,7 +423,7 @@ jobs:
       id-token: write
       packages: read
       pull-requests: read
-    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@test-gb300
     with:
       build_type: pull-request
       package_name: libcudf
@@ -436,13 +436,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_cpp
     with:
       build_type: pull-request
       script: ci/test_cpp.sh
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-python-build:
     needs: [build-details, conda-cpp-build]
     permissions:
@@ -452,7 +450,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@test-gb300
     with:
       build_type: pull-request
       build-datetime: ${{ needs.build-details.outputs.build-datetime }}
@@ -468,7 +466,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@test-gb300
     with:
       build_type: pull-request
       build-datetime: ${{ needs.build-details.outputs.build-datetime }}
@@ -483,13 +481,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_conda && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cudf_polars_nor_dask_cudf
     with:
       build_type: pull-request
       script: "ci/test_python_cudf.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-python-other-tests:
     # Tests for dask_cudf, cudf_polars, custreamz, cudf_kafka are separated for CI parallelism
     needs: [conda-python-build-noarch, changed-files]
@@ -500,15 +496,13 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_conda
     with:
       build_type: pull-request
       # https://github.com/rapidsai/cudf/pull/22381/changes#r3196736965
       container-options: "--cap-add CAP_SYS_PTRACE --shm-size=8g --ulimit=nofile=1000000:1000000"
       script: "ci/test_python_other.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-java-tests:
     needs: [conda-cpp-build, changed-files]
     permissions:
@@ -518,7 +512,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_java
     with:
       build_type: pull-request
@@ -535,7 +529,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_notebooks
     with:
       build_type: pull-request
@@ -552,7 +546,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).build_docs
     with:
       build_type: pull-request
@@ -575,7 +569,7 @@ jobs:
       target-aws-region: ${{ secrets.NVIDIA_DOCS_AWS_REGION }}
       target-aws-secret-access-key: ${{ secrets.NVIDIA_DOCS_AWS_SECRET_ACCESS_KEY }}
       target-s3-bucket: ${{ secrets.NVIDIA_DOCS_S3_BUCKET }}
-    uses: rapidsai/shared-workflows/.github/workflows/publish-api-docs.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/publish-api-docs.yaml@test-gb300
     with:
       docs-projects: cudf,libcudf
       dry-run: true
@@ -589,7 +583,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@test-gb300
     with:
       # build for every combination of arch and CUDA version, but only for the latest Python
       matrix_filter: group_by([.ARCH, (.CUDA_VER|split(".")|map(tonumber)|.[0])]) | map(max_by(.PY_VER|split(".")|map(tonumber)))
@@ -608,7 +602,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@test-gb300
     with:
       # build for every combination of arch and CUDA version, but only for the latest Python
       matrix_filter: group_by([.ARCH, (.CUDA_VER|split(".")|map(tonumber)|.[0])]) | map(max_by(.PY_VER|split(".")|map(tonumber)))
@@ -627,7 +621,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@test-gb300
     with:
       # Build a wheel for each CUDA x ARCH x minimum supported Python version
       matrix_filter: group_by({CUDA_VER, ARCH}) | map(min_by(.PY_VER | split(".") | map(tonumber)))
@@ -646,13 +640,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels
     with:
       build_type: pull-request
       script: ci/test_wheel_cudf_streaming.sh
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   wheel-build-pylibcudf:
     needs: [build-details, checks, wheel-build-libcudf]
     permissions:
@@ -662,7 +654,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@test-gb300
     with:
       build_type: pull-request
       build-datetime: ${{ needs.build-details.outputs.build-datetime }}
@@ -681,7 +673,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@test-gb300
     with:
       build_type: pull-request
       build-datetime: ${{ needs.build-details.outputs.build-datetime }}
@@ -700,13 +692,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cudf_polars_nor_dask_cudf
     with:
       build_type: pull-request
       script: ci/test_wheel_cudf.sh
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   wheel-build-cudf-polars:
     needs: [build-details, wheel-build-pylibcudf]
     permissions:
@@ -716,7 +706,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@test-gb300
     with:
       # This selects "ARCH=amd64 + the latest supported Python + CUDA".
       matrix_filter: map(select(.ARCH == "amd64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
@@ -736,11 +726,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cudf_nor_dask_cudf
     with:
       # This selects "ARCH=amd64 + the latest supported Python + CUDA".
-      matrix_filter: map(select(.ARCH == "amd64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
+      matrix_filter: map(select(.ARCH == "arm64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
       build_type: pull-request
       container-options: "--cap-add CAP_SYS_PTRACE --shm-size=8g --ulimit=nofile=1000000:1000000"
       script: "env POLARS_VERSIONS=endpoints ci/test_wheel_cudf_polars.sh"
@@ -753,11 +743,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cudf_nor_dask_cudf
     with:
       # This selects "ARCH=amd64 + the latest supported Python + CUDA".
-      matrix_filter: map(select(.ARCH == "amd64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
+      matrix_filter: map(select(.ARCH == "arm64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
       build_type: pull-request
       script: "ci/test_cudf_polars_polars_tests.sh"
   wheel-build-dask-cudf:
@@ -769,7 +759,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@test-gb300
     with:
       # This selects "ARCH=amd64 + the latest supported Python + CUDA".
       matrix_filter: map(select(.ARCH == "amd64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
@@ -789,11 +779,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels && fromJSON(needs.changed-files.outputs.changed_file_groups).not_cudf_polars
     with:
       # This selects "ARCH=amd64 + the latest supported Python + CUDA".
-      matrix_filter: map(select(.ARCH == "amd64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
+      matrix_filter: map(select(.ARCH == "arm64")) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
       build_type: pull-request
       script: ci/test_wheel_dask_cudf.sh
   # build cudf with benchmarks but disable for cudf_streaming since we can't use MPI/UCX in wheel-based builds
@@ -806,7 +796,7 @@ jobs:
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
     needs: telemetry-setup
-    uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@test-gb300
     with:
       arch: '["amd64", "arm64"]'
       cuda: '["13.3"]'
@@ -829,12 +819,11 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@test-gb300
     if: (fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels || fromJSON(needs.changed-files.outputs.changed_file_groups).test_cudf_pandas) && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cudf_polars_nor_dask_cudf
     with:
       # This selects the latest supported Python + CUDA minor versions for each ARCH/CUDA major version combo
-      # Filter out GB300 due to https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: group_by([(.ARCH), (.CUDA_VER|split(".")|map(tonumber)|.[0])]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))])) | map(select(.GPU != "gb300" and .GPU != "gh200"))
+      matrix_filter: group_by([(.ARCH), (.CUDA_VER|split(".")|map(tonumber)|.[0])]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
       build_type: pull-request
       script: ci/cudf_pandas_scripts/run_tests.sh
   third-party-integration-tests-cudf-pandas:
@@ -846,7 +835,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_conda && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cpp_nor_cudf_polars_nor_dask_cudf
     with:
       build_type: pull-request
@@ -867,7 +856,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_conda && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cudf_polars_nor_dask_cudf
     with:
       build_type: pull-request
@@ -888,7 +877,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_wheels && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cpp_nor_cudf_polars_nor_dask_cudf
     with:
       build_type: pull-request
@@ -907,7 +896,7 @@ jobs:
       packages: read
       pull-requests: read
     secrets: inherit # zizmor: ignore[secrets-inherit]
-    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@main
+    uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@test-gb300
     if: fromJSON(needs.changed-files.outputs.changed_file_groups).test_python_conda && fromJSON(needs.changed-files.outputs.changed_file_groups).neither_cpp_nor_cudf_polars_nor_dask_cudf
     with:
       build_type: pull-request
diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml
index b15763c2654d..47613a0306ca 100644
--- a/.github/workflows/test.yaml
+++ b/.github/workflows/test.yaml
@@ -54,8 +54,6 @@ jobs:
       date: ${{ inputs.date }}
       script: ci/test_cpp.sh
       sha: ${{ inputs.sha }}
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-cpp-benchmark-tests:
     permissions:
       actions: read
@@ -123,8 +121,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_python_cudf.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-python-other-tests:
     # Tests for dask_cudf, custreamz, cudf_kafka are separated for CI parallelism
     permissions:
@@ -141,8 +137,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_python_other.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-java-tests:
     permissions:
       actions: read
@@ -224,8 +218,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: ci/cudf_pandas_scripts/run_tests.sh
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   third-party-integration-tests-cudf-pandas:
     permissions:
       actions: read
@@ -279,8 +271,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_wheel_cudf_polars.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   cudf-polars-polars-tests:
     permissions:
       actions: read
@@ -296,8 +286,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_cudf_polars_polars_tests.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   narwhals-tests:
     permissions:
       actions: read
diff --git a/cpp/examples/parquet_inspect/parquet_inspect_utils.cpp b/cpp/examples/parquet_inspect/parquet_inspect_utils.cpp
index e8c846aaca02..1ff219ddb965 100644
--- a/cpp/examples/parquet_inspect/parquet_inspect_utils.cpp
+++ b/cpp/examples/parquet_inspect/parquet_inspect_utils.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -8,6 +8,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda_memcpy.hpp>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/io/experimental/hybrid_scan.hpp>
 #include <cudf/io/parquet_io_utils.hpp>
 #include <cudf/io/parquet_schema.hpp>
@@ -120,7 +121,8 @@ auto make_index_column(cudf::size_type num_rows, rmm::cuda_stream_view stream)
 {
   std::vector<cudf::size_type> data(num_rows);
   std::iota(data.begin(), data.end(), 0);
-  auto buffer = rmm::device_buffer(data.data(), num_rows * sizeof(int64_t), stream);
+  auto buffer = cudf::detail::make_device_buffer_async(
+    cudf::host_span<cudf::size_type const>{data}, stream, cudf::get_current_device_resource_ref());
   return std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<cudf::size_type>()},
                                         num_rows,
                                         std::move(buffer),
@@ -140,7 +142,8 @@ auto make_index_column(cudf::size_type num_rows, rmm::cuda_stream_view stream)
 template <typename T>
 auto make_column(cudf::host_span<T const> host_data, rmm::cuda_stream_view stream)
 {
-  auto device_buffer = rmm::device_buffer(host_data.data(), host_data.size() * sizeof(T), stream);
+  auto device_buffer = cudf::detail::make_device_buffer_async(
+    host_data, stream, cudf::get_current_device_resource_ref());
   return std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<T>()},
                                         host_data.size(),
                                         std::move(device_buffer),
@@ -172,8 +175,8 @@ auto make_page_data_list_column(cudf::host_span<T const> data,
 
   auto offsets_column = make_column<cudf::size_type>(col_page_offsets, stream);
 
-  auto page_data_buffer =
-    rmm::device_buffer(data.data(), num_pages_this_column * sizeof(int64_t), stream);
+  auto page_data_buffer = cudf::detail::make_device_buffer_async(
+    data.subspan(0, num_pages_this_column), stream, cudf::get_current_device_resource_ref());
 
   auto page_data_column =
     std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<int64_t>()},
@@ -279,11 +282,17 @@ void write_rowgroup_metadata(cudf::io::parquet::FileMetaData const& metadata,
   columns.emplace_back(make_index_column(num_row_groups, stream));
 
   auto row_offsets_buffer =
-    rmm::device_buffer(row_group_row_offsets.data(), num_row_groups * sizeof(int64_t), stream);
+    cudf::detail::make_device_buffer_async(cudf::host_span<int64_t const>{row_group_row_offsets},
+                                           stream,
+                                           cudf::get_current_device_resource_ref());
   auto row_counts_buffer =
-    rmm::device_buffer(row_group_row_counts.data(), num_row_groups * sizeof(int64_t), stream);
+    cudf::detail::make_device_buffer_async(cudf::host_span<int64_t const>{row_group_row_counts},
+                                           stream,
+                                           cudf::get_current_device_resource_ref());
   auto byte_offsets_buffer =
-    rmm::device_buffer(row_group_byte_offsets.data(), num_row_groups * sizeof(int64_t), stream);
+    cudf::detail::make_device_buffer_async(cudf::host_span<int64_t const>{row_group_byte_offsets},
+                                           stream,
+                                           cudf::get_current_device_resource_ref());
 
   columns.emplace_back(std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<int64_t>()},
                                                       num_row_groups,
diff --git a/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp b/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp
index 3764a1f345b4..784cce8062c7 100644
--- a/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp
+++ b/cpp/include/cudf/detail/utilities/cuda_memcpy.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -19,6 +19,11 @@ enum class host_memory_kind : uint8_t { PINNED, PAGEABLE };
 void cuda_memcpy_async_impl(
   void* dst, void const* src, size_t size, host_memory_kind kind, rmm::cuda_stream_view stream);
 
+[[nodiscard]] cudaError_t memcpy_h2d_async(void* dst,
+                                           void const* src,
+                                           size_t size,
+                                           rmm::cuda_stream_view stream);
+
 /**
  * @brief Wrapper around cudaMemcpyBatchAsync
  *
@@ -70,9 +75,10 @@ void cuda_memcpy_async_impl(
                                        rmm::cuda_stream_view stream);
 
 /**
- * @brief Asynchronously copies data from host to device memory.
+ * @brief Asynchronously copies data from host to device memory
  *
- * Implementation may use different strategies depending on the size and type of host data.
+ * The destination copy remains stream ordered. The host source may be released or changed as soon
+ * as this function returns.
  *
  * @param dst Destination device memory
  * @param src Source host memory
@@ -82,12 +88,7 @@ template <typename T>
 void cuda_memcpy_async(device_span<T> dst, host_span<T const> src, rmm::cuda_stream_view stream)
 {
   CUDF_EXPECTS(dst.size() == src.size(), "Mismatched sizes in cuda_memcpy_async");
-  auto const is_pinned = src.is_device_accessible();
-  cuda_memcpy_async_impl(dst.data(),
-                         src.data(),
-                         src.size_bytes(),
-                         is_pinned ? host_memory_kind::PINNED : host_memory_kind::PAGEABLE,
-                         stream);
+  CUDF_CUDA_TRY(memcpy_h2d_async(dst.data(), src.data(), src.size_bytes(), stream));
 }
 
 /**
diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp
index f82ba7aaece2..7c7b836e8a2b 100644
--- a/cpp/include/cudf/detail/utilities/vector_factories.hpp
+++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -21,6 +21,7 @@
 #include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_buffer.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/resource_ref.hpp>
 
@@ -77,7 +78,8 @@ rmm::device_uvector<T> make_zeroed_device_uvector(std::size_t size,
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
  * `host_span`
  *
- * @note This function does not synchronize `stream`.
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
  *
  * @tparam T The type of the data to copy (may be const-qualified)
  * @param source_data The host_span of data to deep copy
@@ -99,7 +101,8 @@ rmm::device_uvector<std::remove_cv_t<T>> make_device_uvector_async(
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a host
  * container
  *
- * @note This function does not synchronize `stream`.
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
  *
  * @tparam Container The type of the container to copy from
  * @tparam T The type of the data to copy
@@ -119,7 +122,8 @@ rmm::device_uvector<typename Container::value_type> make_device_uvector_async(
 /**
  * @brief Asynchronously construct a `device_uvector` from a `std::vector`
  *
- * @note This function does not synchronize `stream`.
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
  *
  * @tparam T The type of the data to copy
  * @tparam Allocator The allocator type of the std::vector
@@ -136,6 +140,32 @@ rmm::device_uvector<T> make_device_uvector_async(std::vector<T, Allocator> const
   return make_device_uvector_async(host_span<T const>{source_data}, stream, mr);
 }
 
+/**
+ * @brief Asynchronously construct a `device_buffer` containing a deep copy of host data
+ *
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
+ *
+ * @tparam T The type of the data to copy (may be const-qualified)
+ * @param source_data The host data to deep copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_buffer
+ * @return A device_buffer containing the copied data
+ */
+template <typename T>
+rmm::device_buffer make_device_buffer_async(host_span<T> source_data,
+                                            rmm::cuda_stream_view stream,
+                                            rmm::device_async_resource_ref mr)
+{
+  using value_type = std::remove_cv_t<T>;
+  rmm::device_buffer ret(source_data.size_bytes(), stream, mr);
+  cuda_memcpy_async<value_type>(
+    device_span<value_type>{static_cast<value_type*>(ret.data()), source_data.size()},
+    host_span<value_type const>{source_data},
+    stream);
+  return ret;
+}
+
 /**
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
  * `device_span`
diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp
index f0a7a64e3d29..0985f98debd4 100644
--- a/cpp/include/cudf_test/column_wrapper.hpp
+++ b/cpp/include/cudf_test/column_wrapper.hpp
@@ -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
  */
 
@@ -160,8 +160,10 @@ rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
   auto transform_begin = thrust::make_transform_iterator(begin, transformer);
   auto const size      = cudf::distance(begin, end);
   auto const elements  = thrust::host_vector<ElementTo>(transform_begin, transform_begin + size);
-  return rmm::device_buffer{
-    elements.data(), size * sizeof(ElementTo), cudf::test::get_default_stream()};
+  return cudf::detail::make_device_buffer_async(
+    cudf::host_span<ElementTo const>{elements.data(), elements.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
 }
 
 // The two signatures below are identical to the above overload apart from
@@ -190,8 +192,10 @@ rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
   auto transform_begin = thrust::make_transform_iterator(begin, transformer);
   auto const size      = cudf::distance(begin, end);
   auto const elements  = thrust::host_vector<RepType>(transform_begin, transform_begin + size);
-  return rmm::device_buffer{
-    elements.data(), size * sizeof(RepType), cudf::test::get_default_stream()};
+  return cudf::detail::make_device_buffer_async(
+    cudf::host_span<RepType const>{elements.data(), elements.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
 }
 
 /**
@@ -221,8 +225,10 @@ rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
   auto transformer_begin = thrust::make_transform_iterator(begin, to_rep);
   auto const size        = cudf::distance(begin, end);
   auto const elements = thrust::host_vector<RepType>(transformer_begin, transformer_begin + size);
-  return rmm::device_buffer{
-    elements.data(), size * sizeof(RepType), cudf::test::get_default_stream()};
+  return cudf::detail::make_device_buffer_async(
+    cudf::host_span<RepType const>{elements.data(), elements.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
 }
 //! @endcond
 
@@ -277,9 +283,10 @@ std::pair<rmm::device_buffer, cudf::size_type> make_null_mask(ValidityIterator b
                                                               ValidityIterator end)
 {
   auto [null_mask, null_count] = make_null_mask_vector(begin, end);
-  auto d_mask                  = rmm::device_buffer{null_mask.data(),
-                                   cudf::bitmask_allocation_size_bytes(cudf::distance(begin, end)),
-                                   cudf::test::get_default_stream()};
+  auto d_mask                  = cudf::detail::make_device_buffer_async(
+    cudf::host_span<bitmask_type const>{null_mask.data(), null_mask.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
   return {std::move(d_mask), null_count};
 }
 
@@ -563,12 +570,14 @@ class fixed_point_column_wrapper : public detail::column_wrapper {
     auto const id        = type_to_id<numeric::fixed_point<Rep, numeric::Radix::BASE_10>>();
     auto const data_type = cudf::data_type{id, static_cast<int32_t>(scale)};
 
-    wrapped.reset(new cudf::column{
-      data_type,
-      size,
-      rmm::device_buffer{elements.data(), size * sizeof(Rep), cudf::test::get_default_stream()},
-      rmm::device_buffer{},
-      0});
+    wrapped.reset(new cudf::column{data_type,
+                                   size,
+                                   cudf::detail::make_device_buffer_async(
+                                     cudf::host_span<Rep const>{elements.data(), elements.size()},
+                                     cudf::test::get_default_stream(),
+                                     cudf::get_current_device_resource_ref()),
+                                   rmm::device_buffer{},
+                                   0});
   }
 
   /**
@@ -629,12 +638,14 @@ class fixed_point_column_wrapper : public detail::column_wrapper {
     auto const id                = type_to_id<numeric::fixed_point<Rep, numeric::Radix::BASE_10>>();
     auto const data_type         = cudf::data_type{id, static_cast<int32_t>(scale)};
     auto [null_mask, null_count] = detail::make_null_mask(v, v + size);
-    wrapped.reset(new cudf::column{
-      data_type,
-      size,
-      rmm::device_buffer{elements.data(), size * sizeof(Rep), cudf::test::get_default_stream()},
-      std::move(null_mask),
-      null_count});
+    wrapped.reset(new cudf::column{data_type,
+                                   size,
+                                   cudf::detail::make_device_buffer_async(
+                                     cudf::host_span<Rep const>{elements.data(), elements.size()},
+                                     cudf::test::get_default_stream(),
+                                     cudf::get_current_device_resource_ref()),
+                                   std::move(null_mask),
+                                   null_count});
   }
 
   /**
diff --git a/cpp/libcudf_streaming/tests/streaming/test_channel_metadata.cpp b/cpp/libcudf_streaming/tests/streaming/test_channel_metadata.cpp
index 881cfd488b6b..5faf98a1df81 100644
--- a/cpp/libcudf_streaming/tests/streaming/test_channel_metadata.cpp
+++ b/cpp/libcudf_streaming/tests/streaming/test_channel_metadata.cpp
@@ -6,6 +6,7 @@
 #include <cudf_test/cudf_gtest.hpp>
 
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/utilities/default_stream.hpp>
 
@@ -166,7 +167,8 @@ class StreamingChannelMetadataGPU : public ::testing::Test {
 
   std::shared_ptr<table_chunk> make_chunk(std::vector<int32_t> vals)
   {
-    rmm::device_buffer buf(vals.data(), vals.size() * sizeof(int32_t), stream);
+    auto buf = cudf::detail::make_device_buffer_async(
+      cudf::host_span<int32_t const>{vals}, stream, cudf::get_current_device_resource_ref());
     auto col = std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::INT32},
                                               static_cast<cudf::size_type>(vals.size()),
                                               std::move(buf),
diff --git a/cpp/src/interop/from_arrow_host.cu b/cpp/src/interop/from_arrow_host.cu
index 8f6cdf4483e2..068f99a99b00 100644
--- a/cpp/src/interop/from_arrow_host.cu
+++ b/cpp/src/interop/from_arrow_host.cu
@@ -110,7 +110,8 @@ std::pair<std::unique_ptr<rmm::device_buffer>, size_type> get_mask_buffer(
   auto const copy_size    = cudf::util::div_rounding_up_safe(num_rows + bit_index, bits_in_byte);
 
   auto mask = rmm::device_uvector<bitmask_type>(padded_words, stream, mr);
-  CUDF_CUDA_TRY(cudf::detail::memcpy_async(mask.data(), bitmap + offset_index, copy_size, stream));
+  CUDF_CUDA_TRY(
+    cudf::detail::memcpy_h2d_async(mask.data(), bitmap + offset_index, copy_size, stream));
 
   if (mask_words > 0 && bit_index > 0) {
     auto dest_mask = rmm::device_uvector<bitmask_type>(padded_words, stream, mr);
@@ -158,10 +159,10 @@ struct dispatch_copy_from_arrow_host {
 
     auto col = make_fixed_width_column(type, num_rows, mask_state::UNALLOCATED, stream, mr);
     auto mutable_column_view = col->mutable_view();
-    CUDF_CUDA_TRY(cudf::detail::memcpy_async(mutable_column_view.data<DeviceType>(),
-                                             data_buffer + input->offset,
-                                             sizeof(DeviceType) * num_rows,
-                                             stream));
+    CUDF_CUDA_TRY(cudf::detail::memcpy_h2d_async(mutable_column_view.data<DeviceType>(),
+                                                 data_buffer + input->offset,
+                                                 sizeof(DeviceType) * num_rows,
+                                                 stream));
 
     if (!skip_mask) {
       auto [mask, null_count] = get_mask_buffer(input, stream, mr);
@@ -190,7 +191,7 @@ std::unique_ptr<column> dispatch_copy_from_arrow_host::operator()<bool>(ArrowSch
 
   auto data = rmm::device_uvector<bitmask_type>(data_words, stream, mr);
   CUDF_CUDA_TRY(
-    cudf::detail::memcpy_async(data.data(), data_buffer + offset_index, copy_size, stream));
+    cudf::detail::memcpy_h2d_async(data.data(), data_buffer + offset_index, copy_size, stream));
 
   if (data_words > 0 && bit_index > 0) {
     auto dest_data = rmm::device_uvector<bitmask_type>(data_words, stream, mr);
diff --git a/cpp/src/interop/from_arrow_host_strings.cu b/cpp/src/interop/from_arrow_host_strings.cu
index 54f957d5522d..80cb71029d18 100644
--- a/cpp/src/interop/from_arrow_host_strings.cu
+++ b/cpp/src/interop/from_arrow_host_strings.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -53,7 +53,7 @@ std::unique_ptr<column> from_arrow_string(ArrowSchemaView const* schema,
 
   rmm::device_buffer chars(char_data_length, stream, mr);
   auto const* chars_data = static_cast<uint8_t const*>(input->buffers[chars_buffer_idx]) + offset;
-  CUDF_CUDA_TRY(cudf::detail::memcpy_async(chars.data(), chars_data, chars.size(), stream));
+  CUDF_CUDA_TRY(cudf::detail::memcpy_h2d_async(chars.data(), chars_data, chars.size(), stream));
 
   return make_strings_column(static_cast<size_type>(input->length),
                              std::move(offsets_column),
@@ -77,14 +77,16 @@ std::unique_ptr<column> from_arrow_stringview(ArrowSchemaView const* schema,
   // first copy stringview array to device
   auto items   = view.buffer_views[stringview_vector_idx].data.as_binary_view;
   auto d_items = rmm::device_uvector<ArrowBinaryView>(input->length, stream, mr);
-  CUDF_CUDA_TRY(cudf::detail::memcpy_async(
+  CUDF_CUDA_TRY(cudf::detail::memcpy_h2d_async(
     d_items.data(), items + input->offset, input->length * sizeof(ArrowBinaryView), stream));
 
   // then copy variadic buffers to device
   auto variadics     = std::vector<rmm::device_buffer>();
   auto variadic_ptrs = std::vector<char const*>();
   for (auto i = 0L; i < view.n_variadic_buffers; ++i) {
-    variadics.emplace_back(view.variadic_buffers[i], view.variadic_buffer_sizes[i], stream);
+    variadics.emplace_back(view.variadic_buffer_sizes[i], stream, mr);
+    CUDF_CUDA_TRY(cudf::detail::memcpy_h2d_async(
+      variadics.back().data(), view.variadic_buffers[i], view.variadic_buffer_sizes[i], stream));
     variadic_ptrs.push_back(static_cast<char const*>(variadics.back().data()));
   }
 
diff --git a/cpp/src/interop/from_arrow_stream.cu b/cpp/src/interop/from_arrow_stream.cu
index c6825d503dc5..c23738ad422e 100644
--- a/cpp/src/interop/from_arrow_stream.cu
+++ b/cpp/src/interop/from_arrow_stream.cu
@@ -78,16 +78,11 @@ std::unique_ptr<table> from_arrow_stream(ArrowArrayStream* input,
   NANOARROW_THROW_NOT_OK(ArrowArrayStreamGetSchema(input, &schema, nullptr));
 
   std::vector<std::unique_ptr<cudf::table>> chunks;
-  // Keep each input chunk alive until the stream has executed the host-to-device copies enqueued
-  // by `from_arrow`. Those copies use `cudaMemcpyBatchAsync` with `cudaMemcpySrcAccessOrderStream`,
-  // which defers reading the host source until the stream reaches the copy.
-  std::vector<nanoarrow::UniqueArray> sources;
   while (true) {
     nanoarrow::UniqueArray chunk;
     NANOARROW_THROW_NOT_OK(ArrowArrayStreamGetNext(input, chunk.get(), nullptr));
     if (chunk->release == nullptr) { break; }
-    sources.push_back(std::move(chunk));
-    chunks.push_back(from_arrow(&schema, sources.back().get(), stream, mr));
+    chunks.push_back(from_arrow(&schema, chunk.get(), stream, mr));
   }
   input->release(input);
 
@@ -139,17 +134,11 @@ std::unique_ptr<column> from_arrow_stream_column(ArrowArrayStream* input,
   NANOARROW_THROW_NOT_OK(ArrowArrayStreamGetSchema(input, &schema, nullptr));
 
   std::vector<std::unique_ptr<cudf::column>> chunks;
-  // Keep each input chunk alive until the stream has executed the host-to-device copies enqueued
-  // by `from_arrow_column`. Those copies use `cudaMemcpyBatchAsync` with
-  // `cudaMemcpySrcAccessOrderStream`, which defers reading the host source until the stream reaches
-  // the copy.
-  std::vector<nanoarrow::UniqueArray> sources;
   while (true) {
     nanoarrow::UniqueArray chunk;
     NANOARROW_THROW_NOT_OK(ArrowArrayStreamGetNext(input, chunk.get(), nullptr));
     if (chunk->release == nullptr) { break; }
-    sources.push_back(std::move(chunk));
-    chunks.push_back(from_arrow_column(&schema, sources.back().get(), stream, mr));
+    chunks.push_back(from_arrow_column(&schema, chunk.get(), stream, mr));
   }
   input->release(input);
 
diff --git a/cpp/src/io/json/nested_json_gpu.cu b/cpp/src/io/json/nested_json_gpu.cu
index 2d35da95baba..c24183d80094 100644
--- a/cpp/src/io/json/nested_json_gpu.cu
+++ b/cpp/src/io/json/nested_json_gpu.cu
@@ -2106,10 +2106,8 @@ std::pair<std::unique_ptr<column>, std::vector<column_name_info>> json_column_to
     [stream, mr](json_column const& json_col) -> std::pair<rmm::device_buffer, size_type> {
     auto const null_count = json_col.current_offset - json_col.valid_count;
     if (null_count == 0) { return {rmm::device_buffer{}, null_count}; }
-    return {rmm::device_buffer{json_col.validity.data(),
-                               bitmask_allocation_size_bytes(json_col.current_offset),
-                               stream,
-                               mr},
+    return {cudf::detail::make_device_buffer_async(
+              cudf::host_span<bitmask_type const>{json_col.validity}, stream, mr),
             null_count};
   };
 
diff --git a/cpp/src/io/parquet/experimental/deletion_vectors_helpers.cu b/cpp/src/io/parquet/experimental/deletion_vectors_helpers.cu
index 00cb25bc28b1..bc2c36818d13 100644
--- a/cpp/src/io/parquet/experimental/deletion_vectors_helpers.cu
+++ b/cpp/src/io/parquet/experimental/deletion_vectors_helpers.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
diff --git a/cpp/src/io/utilities/hostdevice_vector.hpp b/cpp/src/io/utilities/hostdevice_vector.hpp
index 20aa81ec3455..70c31a1f279d 100644
--- a/cpp/src/io/utilities/hostdevice_vector.hpp
+++ b/cpp/src/io/utilities/hostdevice_vector.hpp
@@ -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
  */
 
diff --git a/cpp/src/strings/combine/join.cu b/cpp/src/strings/combine/join.cu
index aadd3cc83752..92c50384c078 100644
--- a/cpp/src/strings/combine/join.cu
+++ b/cpp/src/strings/combine/join.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
  */
 
diff --git a/cpp/src/utilities/cuda_memcpy.cu b/cpp/src/utilities/cuda_memcpy.cu
index 8e961a9cf633..d505c301bfea 100644
--- a/cpp/src/utilities/cuda_memcpy.cu
+++ b/cpp/src/utilities/cuda_memcpy.cu
@@ -51,16 +51,17 @@ void copy_pageable(void* dst, void const* src, std::size_t size, rmm::cuda_strea
   CUDF_CUDA_TRY(cudf::detail::memcpy_async(dst, src, size, stream));
 }
 
-};  // namespace
-
-cudaError_t memcpy_batch_async(void* const* dsts,
-                               void const* const* srcs,
-                               std::size_t const* sizes,
-                               std::size_t count,
-                               rmm::cuda_stream_view stream)
+enum class host_source_access_order : uint8_t { STREAM, DURING_API_CALL };
+
+cudaError_t memcpy_batch_async_impl(void* const* dsts,
+                                    void const* const* srcs,
+                                    std::size_t const* sizes,
+                                    std::size_t count,
+                                    rmm::cuda_stream_view stream,
+                                    host_source_access_order source_access_order)
 {
-// Uses cudaMemcpyBatchAsync for CUDA 13.0+ to avoid driver-side locking overhead.
-// cudaMemcpyBatchAsync does not support the default stream.
+  // Uses cudaMemcpyBatchAsync for CUDA 13.0+ to avoid driver-side locking overhead.
+  // cudaMemcpyBatchAsync does not support the default stream.
 #if CUDART_VERSION >= 13000
   if (!stream.is_default()) {
     // Filter out invalid copies (nullptr dst/src or size==0);
@@ -90,7 +91,11 @@ cudaError_t memcpy_batch_async(void* const* dsts,
       count = valid_dsts.size();
     }
 
-    cudaMemcpyAttributes attrs = {.srcAccessOrder = cudaMemcpySrcAccessOrderStream,
+    auto const cuda_source_access_order =
+      source_access_order == host_source_access_order::DURING_API_CALL
+        ? cudaMemcpySrcAccessOrderDuringApiCall
+        : cudaMemcpySrcAccessOrderStream;
+    cudaMemcpyAttributes attrs = {.srcAccessOrder = cuda_source_access_order,
                                   .flags          = cudaMemcpyFlagPreferOverlapWithCompute};
     std::size_t attrs_idxs     = 0;
     return cudaMemcpyBatchAsync(dsts, srcs, sizes, count, &attrs, &attrs_idxs, 1, stream.value());
@@ -101,7 +106,21 @@ cudaError_t memcpy_batch_async(void* const* dsts,
       cudaMemcpyAsync(dsts[i], srcs[i], sizes[i], cudaMemcpyDefault, stream.value());
     if (status != cudaSuccess) { return status; }
   }
-  return cudaSuccess;
+  return source_access_order == host_source_access_order::DURING_API_CALL
+           ? cudaStreamSynchronize(stream.value())
+           : cudaSuccess;
+}
+
+}  // namespace
+
+cudaError_t memcpy_batch_async(void* const* dsts,
+                               void const* const* srcs,
+                               std::size_t const* sizes,
+                               std::size_t count,
+                               rmm::cuda_stream_view stream)
+{
+  return memcpy_batch_async_impl(
+    dsts, srcs, sizes, count, stream, host_source_access_order::STREAM);
 }
 
 cudaError_t memcpy_async(void* dst, void const* src, size_t count, rmm::cuda_stream_view stream)
@@ -113,6 +132,17 @@ cudaError_t memcpy_async(void* dst, void const* src, size_t count, rmm::cuda_str
   return memcpy_batch_async(&dst, &src, &count, 1, stream);
 }
 
+CUDF_EXPORT cudaError_t memcpy_h2d_async(void* dst,
+                                         void const* src,
+                                         size_t count,
+                                         rmm::cuda_stream_view stream)
+{
+  if (count == 0) { return cudaSuccess; }
+
+  return memcpy_batch_async_impl(
+    &dst, &src, &count, 1, stream, host_source_access_order::DURING_API_CALL);
+}
+
 void cuda_memcpy_async_impl(
   void* dst, void const* src, size_t size, host_memory_kind kind, rmm::cuda_stream_view stream)
 {
diff --git a/cpp/tests/io/experimental/hybrid_scan_common.cpp b/cpp/tests/io/experimental/hybrid_scan_common.cpp
index 339009ecf7e5..000b3abd28cb 100644
--- a/cpp/tests/io/experimental/hybrid_scan_common.cpp
+++ b/cpp/tests/io/experimental/hybrid_scan_common.cpp
@@ -13,6 +13,7 @@
 #include <cudf/concatenate.hpp>
 #include <cudf/copying.hpp>
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/io/parquet_io_utils.hpp>
 #include <cudf/table/table_view.hpp>
 #include <cudf/utilities/error.hpp>
@@ -345,8 +346,10 @@ std::pair<std::unique_ptr<cudf::table>, std::vector<char>> create_parquet_with_s
 
     auto const make_null_mask = [stream](auto begin, auto end) {
       auto [null_mask, null_count] = cudf::test::detail::make_null_mask_vector(begin, end);
-      auto d_mask                  = rmm::device_buffer{
-        null_mask.data(), cudf::bitmask_allocation_size_bytes(cudf::distance(begin, end)), stream};
+      auto d_mask =
+        cudf::detail::make_device_buffer_async(cudf::host_span<cudf::bitmask_type const>{null_mask},
+                                               stream,
+                                               cudf::get_current_device_resource_ref());
       return std::pair{std::move(d_mask), null_count};
     };
 
diff --git a/cpp/tests/utilities/identify_stream_usage.cpp b/cpp/tests/utilities/identify_stream_usage.cpp
index 00249b830346..c80650cdf22f 100644
--- a/cpp/tests/utilities/identify_stream_usage.cpp
+++ b/cpp/tests/utilities/identify_stream_usage.cpp
@@ -218,6 +218,10 @@ void sanitizer_subscriber::callback(Sanitizer_CallbackDomain domain,
           CHECK_STREAM_ARG(cudaMemcpy3DPeerAsync_ptsz, 7000, stream);
           CHECK_STREAM_ARG(cudaMemcpyAsync, 3020, stream);
           CHECK_STREAM_ARG(cudaMemcpyAsync_ptsz, 7000, stream);
+#if CUDART_VERSION >= 13000
+          CHECK_STREAM_ARG(cudaMemcpyBatchAsync, 13000, stream);
+          CHECK_STREAM_ARG(cudaMemcpyBatchAsync_ptsz, 13000, stream);
+#endif
           CHECK_STREAM_ARG(cudaMemcpyFromSymbolAsync, 3020, stream);
           CHECK_STREAM_ARG(cudaMemcpyFromSymbolAsync_ptsz, 7000, stream);
           CHECK_STREAM_ARG(cudaMemcpyToSymbolAsync, 3020, stream);
diff --git a/cpp/tests/utilities_tests/batched_memcpy_tests.cu b/cpp/tests/utilities_tests/batched_memcpy_tests.cu
index 22c6360b1d3a..5a5b73ffc77c 100644
--- a/cpp/tests/utilities_tests/batched_memcpy_tests.cu
+++ b/cpp/tests/utilities_tests/batched_memcpy_tests.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -13,6 +13,7 @@
 #include <cudf/utilities/memory_resource.hpp>
 #include <cudf/utilities/span.hpp>
 
+#include <rmm/cuda_stream.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/device_vector.hpp>
 
@@ -125,3 +126,43 @@ TEST(BatchedMemcpyTest, BasicTest)
   EXPECT_TRUE(
     std::equal(expected_buffer.begin(), expected_buffer.end(), result_dst_buffer.begin()));
 }
+
+void expect_uvector_source_can_be_modified(rmm::cuda_stream_view stream)
+{
+  auto const mr = cudf::get_current_device_resource_ref();
+  auto source   = std::vector<int32_t>(1024, 1);
+
+  auto result = cudf::detail::make_device_uvector_async(source, stream, mr);
+  std::fill(source.begin(), source.end(), 2);
+
+  auto const actual = cudf::detail::make_std_vector<int32_t>(result, stream);
+  EXPECT_TRUE(std::all_of(actual.cbegin(), actual.cend(), [](auto value) { return value == 1; }));
+}
+
+TEST(BatchedMemcpyTest, UvectorHostSourceCanBeModifiedAfterReturn)
+{
+  rmm::cuda_stream stream;
+  expect_uvector_source_can_be_modified(stream.view());
+}
+
+TEST(BatchedMemcpyTest, UvectorHostSourceCanBeModifiedAfterReturnDefaultStream)
+{
+  expect_uvector_source_can_be_modified(cudf::get_default_stream());
+}
+
+TEST(BatchedMemcpyTest, DeviceBufferHostSourceCanBeModifiedAfterReturn)
+{
+  rmm::cuda_stream stream;
+  auto const stream_view = stream.view();
+  auto const mr          = cudf::get_current_device_resource_ref();
+  auto source            = std::vector<int32_t>(1024, 1);
+
+  auto result =
+    cudf::detail::make_device_buffer_async(cudf::host_span<int32_t const>{source}, stream_view, mr);
+  std::fill(source.begin(), source.end(), 2);
+
+  auto const actual = cudf::detail::make_std_vector<int32_t>(
+    cudf::device_span<int32_t const>{static_cast<int32_t const*>(result.data()), source.size()},
+    stream_view);
+  EXPECT_TRUE(std::all_of(actual.cbegin(), actual.cend(), [](auto value) { return value == 1; }));
+}