Skip to content

Commit

Permalink
Fix host-to-device copy missing sync in strings/duration convert (#17149
Browse files Browse the repository at this point in the history
)

Fixes a missing stream sync when copying a temporary host vector to device. The host vector could be destroyed before the copy is completed. Updates the code to use vector factory function `make_device_uvector_sync()` instead of `cudaMemcpyAsync`

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - Yunsong Wang (https://github.com/PointKernel)

URL: #17149
  • Loading branch information
davidwendt authored Oct 25, 2024
1 parent 7115f20 commit 03777f6
Showing 1 changed file with 3 additions and 6 deletions.
9 changes: 3 additions & 6 deletions cpp/src/strings/convert/convert_durations.cu
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@
#include <cudf/column/column_device_view.cuh>
#include <cudf/detail/null_mask.hpp>
#include <cudf/detail/nvtx/ranges.hpp>
#include <cudf/detail/utilities/vector_factories.hpp>
#include <cudf/strings/convert/convert_durations.hpp>
#include <cudf/strings/detail/convert/int_to_string.cuh>
#include <cudf/strings/detail/strings_children.cuh>
Expand Down Expand Up @@ -152,12 +153,8 @@ struct format_compiler {
}

// create program in device memory
d_items.resize(items.size(), stream);
CUDF_CUDA_TRY(cudaMemcpyAsync(d_items.data(),
items.data(),
items.size() * sizeof(items[0]),
cudaMemcpyDefault,
stream.value()));
d_items = cudf::detail::make_device_uvector_sync(
items, stream, cudf::get_current_device_resource_ref());
}

format_item const* compiled_format_items() { return d_items.data(); }
Expand Down

0 comments on commit 03777f6

Please sign in to comment.