diff --git a/cpp/include/cudf/scalar/scalar.hpp b/cpp/include/cudf/scalar/scalar.hpp index 13a3da14cce..39bd2984095 100644 --- a/cpp/include/cudf/scalar/scalar.hpp +++ b/cpp/include/cudf/scalar/scalar.hpp @@ -706,13 +706,16 @@ class list_scalar : public scalar { */ class struct_scalar : public scalar { public: - struct_scalar() = delete; - ~struct_scalar() = default; - struct_scalar(struct_scalar&& other) = default; - struct_scalar(struct_scalar const& other) = default; + struct_scalar() = delete; + ~struct_scalar() = default; + struct_scalar(struct_scalar&& other) = default; struct_scalar& operator=(struct_scalar const& other) = delete; struct_scalar& operator=(struct_scalar&& other) = delete; + struct_scalar(struct_scalar const& other, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + /** * @brief Construct a new struct scalar object from table_view. * diff --git a/cpp/src/scalar/scalar.cpp b/cpp/src/scalar/scalar.cpp index 045bfbe0327..a8d05e98034 100644 --- a/cpp/src/scalar/scalar.cpp +++ b/cpp/src/scalar/scalar.cpp @@ -520,6 +520,13 @@ list_scalar::list_scalar(list_scalar const& other, column_view list_scalar::view() const { return _data.view(); } +struct_scalar::struct_scalar(struct_scalar const& other, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + : scalar{other, stream, mr}, _data(other._data, stream, mr) +{ +} + struct_scalar::struct_scalar(table_view const& data, bool is_valid, rmm::cuda_stream_view stream, diff --git a/cpp/src/strings/copying/concatenate.cu b/cpp/src/strings/copying/concatenate.cu index 866ff1adbc6..8d77c7da4cc 100644 --- a/cpp/src/strings/copying/concatenate.cu +++ b/cpp/src/strings/copying/concatenate.cu @@ -16,8 +16,7 @@ #include #include -#include -#include +#include #include #include #include @@ -27,12 +26,11 @@ #include #include +#include #include -#include "thrust/iterator/transform_iterator.h" #include -#include -#include +#include #include namespace cudf { @@ -287,12 +285,15 @@ std::unique_ptr concatenate(host_span columns, column_view offsets_child = column->child(strings_column_view::offsets_column_index); column_view chars_child = column->child(strings_column_view::chars_column_index); - auto d_offsets = offsets_child.data() + column_offset; - int32_t bytes_offset = thrust::device_pointer_cast(d_offsets)[0]; + auto bytes_offset = + cudf::detail::get_value(offsets_child, column_offset, stream); // copy the chars column data - auto d_chars = chars_child.data() + bytes_offset; - size_type bytes = thrust::device_pointer_cast(d_offsets)[column_size] - bytes_offset; + auto d_chars = chars_child.data() + bytes_offset; + auto const bytes = + cudf::detail::get_value(offsets_child, column_size + column_offset, stream) - + bytes_offset; + CUDA_TRY( cudaMemcpyAsync(d_new_chars, d_chars, bytes, cudaMemcpyDeviceToDevice, stream.value())); diff --git a/cpp/tests/scalar/scalar_test.cpp b/cpp/tests/scalar/scalar_test.cpp index 2047d815867..b54594fd1c4 100644 --- a/cpp/tests/scalar/scalar_test.cpp +++ b/cpp/tests/scalar/scalar_test.cpp @@ -14,19 +14,12 @@ * limitations under the License. */ -#include -#include -#include #include #include -#include #include -#include #include -#include -#include -#include +#include template struct TypedScalarTest : public cudf::test::BaseFixture {