diff --git a/docs/data-management.md b/docs/data-management.md index 48419a0..5407fec 100644 --- a/docs/data-management.md +++ b/docs/data-management.md @@ -72,16 +72,20 @@ Data representations are thin wrappers -- they hold the data but delegate storag ### GPU Table Representation -**File**: `include/cucascade/data/gpu_data_representation.hpp` +**File**: `include/cucascade/cudf/gpu_data_representation.hpp` -Wraps a `cudf::table` residing in GPU device memory: +Wraps a `cudf::table` residing in GPU device memory. The table is held through an +`owning_table_view` (`include/cucascade/cudf/owning_table_view.hpp`) — a handle that is either an +owned `cudf::table`, a `cudf::table_view` backed by a type-erased external owner, or empty: ```cpp class gpu_table_representation : public idata_representation { - std::unique_ptr _table; + std::size_t _alloc_size; // estimate for view states; actual once materialized + owning_table_view _table; public: - std::unique_ptr release_table(rmm::cuda_stream_view stream); // Transfer ownership + void materialize_table(rmm::cuda_stream_view stream); // View -> owned table, in place + std::unique_ptr release_table(rmm::cuda_stream_view stream); // Transfer ownership std::size_t get_size_in_bytes() const override; std::unique_ptr clone(rmm::cuda_stream_view stream) override; @@ -89,6 +93,9 @@ public: ``` - `clone()` performs a deep copy using `cudf::table(table.view(), stream)` +- `materialize_table()` / `release_table()` realize a view state into an owned table: zero-copy + (column buffers moved out) for exclusively-owned `no_alloc_materializable` owners, otherwise a + copy allocated from the representation's memory space - Owns the `cudf::table` object but not the underlying GPU memory (managed by the allocator) ### Host Table Representation (Direct Copy) diff --git a/include/cucascade/cudf/gpu_data_representation.hpp b/include/cucascade/cudf/gpu_data_representation.hpp index b6d463a..1f58bb8 100644 --- a/include/cucascade/cudf/gpu_data_representation.hpp +++ b/include/cucascade/cudf/gpu_data_representation.hpp @@ -17,6 +17,7 @@ #pragma once +#include #include #include @@ -27,10 +28,11 @@ #include -#include +#include #include #include -#include +#include +#include namespace cucascade { @@ -77,12 +79,17 @@ class gpu_table_representation : public idata_representation { * * @param table_view View of the cuDF table (data ownership lives in @p owner) * @tparam Owner The type of the owner of the cuDF table (e.g., a specific operator or component) - * @param owner Owner of the underlying data (transferred via std::any storage) + * @param owner Owner of the underlying data (type-erased into an owning_table_view; move-only + * owners are supported). Exclusively-owned owners satisfying + * detail::no_alloc_materializable enable zero-copy materialize_table()/ + * release_table(); any other owner is materialized by copy. * @param alloc_size Allocation size in bytes for the data * @param memory_space The memory space where the GPU table resides * @param writer_stream The stream on which @p table_view's data was last written. */ template + requires(!std::same_as, owning_table_view> && + !std::same_as, std::unique_ptr>) gpu_table_representation(cudf::table_view table_view, Owner&& owner, std::size_t alloc_size, @@ -138,13 +145,42 @@ class gpu_table_representation : public idata_representation { /** * @brief Release ownership of the underlying cuDF table * - * After calling this method, this representation no longer owns the table. + * After calling this method, this representation no longer owns the table: + * get_size_in_bytes() returns 0 and get_table_view() returns an empty view. + * A view-state representation is materialized first (zero-copy for + * no-alloc-materializable owners, otherwise copied via this memory space's + * default allocator). Returns nullptr if the table was already released. + * + * STREAM-LINEAGE: @p stream first waits on the recorded writer event (if any) + * so the materializing read is ordered after the producing writes. * * @param stream CUDA stream (used to materialize the table from a view path before release) * @return std::unique_ptr The cuDF table */ std::unique_ptr release_table(rmm::cuda_stream_view stream); + /** + * @brief Materialize a view-state representation into an owned cuDF table in place. + * + * No-op when the representation already owns its table or is empty. For a + * view state, realizes the exposed view into an owned table: zero-copy + * (buffers moved out of the owner) when the owner is exclusively owned and + * no-alloc-materializable, otherwise a copy allocated from this memory + * space's default allocator. get_size_in_bytes() is updated from the + * caller-provided estimate to the materialized table's actual allocation + * size. + * + * STREAM-LINEAGE: @p stream first waits on the recorded writer event (if + * any), then the writer event is re-recorded on @p stream after the + * materialization — wait-then-record keeps the event lineage correct for + * both the move and copy paths. When no writer event was recorded (legacy + * paths), the caller must pass a stream already ordered after the producing + * writes, as with clone(). + * + * @param stream CUDA stream for the materializing copy (and event ordering) + */ + void materialize_table(rmm::cuda_stream_view stream); + /** * @brief Rebind the owned table's device buffers to use @p stream for future deallocation. * @@ -153,7 +189,7 @@ class gpu_table_representation : public idata_representation { * stream they were originally allocated on. No device memory is copied and no kernels are * launched. * - * No-op when this representation holds an owning_table_view: that alternative references + * No-op when this representation holds a view state (not materialized): a view references * memory owned by an external (type-erased) owner, which is responsible for its own * deallocation stream. * @@ -198,14 +234,16 @@ class gpu_table_representation : public idata_representation { [[nodiscard]] cudaEvent_t get_writer_event() const override; private: - struct owning_table_view { - std::any owner; ///< The owner of the cuDF table - std::size_t alloc_size{0}; - cudf::table_view view; ///< A view of the owned table for easy access - }; + /// Bytes attributed to this representation. Captured from the owned table's + /// alloc_size() (owned ctor) or the caller-provided estimate (view ctor); + /// updated to the actual table size by materialize_table(); zeroed by + /// release_table(). Declared before _table so the owned ctor can read the + /// table's size before moving it into _table. + std::size_t _alloc_size{0}; - std::variant, owning_table_view> - _table; ///< cudf::table is the underlying representation of the data + /// The table handle: an owned cudf::table, a type-erased owner + column + /// selection (view state), or empty (after release_table()). + owning_table_view _table; /// Lazily-created CUDA event recording the completion of the most recent /// writer-stream work that produced this representation. Null until the first @@ -214,14 +252,16 @@ class gpu_table_representation : public idata_representation { }; template + requires(!std::same_as, owning_table_view> && + !std::same_as, std::unique_ptr>) gpu_table_representation::gpu_table_representation(cudf::table_view table_view, Owner&& owner, std::size_t alloc_size, cucascade::memory::memory_space& memory_space, rmm::cuda_stream_view writer_stream) : idata_representation(memory_space), - _table( - owning_table_view{std::make_any(std::forward(owner)), alloc_size, table_view}) + _alloc_size(alloc_size), + _table(std::forward(owner), table_view) { // STREAM-LINEAGE: record writer event so cross-stream/cross-device readers // can establish ordering via cudaStreamWaitEvent. diff --git a/include/cucascade/cudf/owning_table_view.hpp b/include/cucascade/cudf/owning_table_view.hpp new file mode 100644 index 0000000..343b155 --- /dev/null +++ b/include/cucascade/cudf/owning_table_view.hpp @@ -0,0 +1,400 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cucascade { + +namespace detail { + +/** + * @brief An exclusively-owned owner whose pointee can surrender its columns as + * owned @c std::unique_ptr objects (the canonical case is + * @c std::unique_ptr). + * + * When this holds, a view built over the owner can be materialized into a + * @c cudf::table by *moving* the surviving column buffers out of the owner + * instead of copying them — no new device buffers are allocated. Owners that + * do not satisfy this concept fall back to a copying materialization. + * + * Copy-constructible owners (e.g. @c std::shared_ptr) are + * deliberately excluded: moving columns out of a shared owner would gut every + * co-owner. The pointee must also expose @c view() so materialization can + * verify the base view actually names the owner's columns before moving them + * out (see the identity guard in @ref owner_backed_view). + */ +template +concept no_alloc_materializable = + !std::copy_constructible && requires(Owner& owner) { + { (*owner).release() } -> std::same_as>>; + { (*owner).view() } -> std::same_as; + }; + +/** + * @brief Whether two column views name the exact same device buffers (recursing + * into children), i.e. moving the underlying column is equivalent to + * keeping the view. + */ +[[nodiscard]] inline bool shallow_identical(cudf::column_view const& lhs, + cudf::column_view const& rhs) +{ + if (lhs.type() != rhs.type() || lhs.size() != rhs.size() || lhs.offset() != rhs.offset() || + lhs.head() != rhs.head() || lhs.null_mask() != rhs.null_mask() || + lhs.num_children() != rhs.num_children()) { + return false; + } + for (cudf::size_type i = 0; i < lhs.num_children(); ++i) { + if (!shallow_identical(lhs.child(i), rhs.child(i))) { return false; } + } + return true; +} + +/// @copydoc shallow_identical(cudf::column_view const&, cudf::column_view const&) +[[nodiscard]] inline bool shallow_identical(cudf::table_view const& lhs, + cudf::table_view const& rhs) +{ + if (lhs.num_columns() != rhs.num_columns() || lhs.num_rows() != rhs.num_rows()) { return false; } + for (cudf::size_type i = 0; i < lhs.num_columns(); ++i) { + if (!shallow_identical(lhs.column(i), rhs.column(i))) { return false; } + } + return true; +} + +//===----------------------------------------------------------------------===// +// owner_backed_view +//===----------------------------------------------------------------------===// +/** + * @brief A type-erased data owner paired with a column selection describing the + * currently-exposed @c cudf::table_view. + * + * The selection is stored as indices into the owner's *original* columns so + * that @ref reorder_columns / @ref drop_columns are pure index manipulations + * (no allocation) and so that @ref materialize can, when the owner supports it + * (@ref no_alloc_materializable), move the surviving column buffers out of the + * owner rather than copying them. + */ +class owner_backed_view { + public: + /// Type-erase @p owner (kept alive for the view's lifetime) together with the + /// @p base_view that names its columns. The initial selection exposes every + /// column of @p base_view in order. + template + owner_backed_view(Owner&& owner, cudf::table_view base_view) + : _model( + std::make_unique>>(std::forward(owner), base_view)), + _selection(identity_selection(base_view.num_columns())) + { + } + + [[nodiscard]] cudf::table_view view() const { return _model->full_view().select(_selection); } + + [[nodiscard]] std::size_t n_columns() const { return _selection.size(); } + + /// Row count of the underlying data, independent of which columns are + /// currently exposed. + [[nodiscard]] cudf::size_type num_rows() const { return _model->full_view().num_rows(); } + + /// Swap the columns at the given current positions, applying each swap in + /// order. Pure index manipulation; allocates nothing. Throws + /// @c std::out_of_range if any position is outside the current view. + void reorder_columns(std::span> swaps) + { + for (auto const& [a, b] : swaps) { + check_position(a); + check_position(b); + std::swap(_selection[a], _selection[b]); + } + } + + /// Drop the columns at the given current positions. Pure index manipulation; + /// allocates nothing. Throws @c std::out_of_range if any position is outside + /// the current view. + void drop_columns(std::span positions) + { + std::vector sorted(positions.begin(), positions.end()); + std::sort(sorted.begin(), sorted.end(), std::greater<>{}); + sorted.erase(std::unique(sorted.begin(), sorted.end()), sorted.end()); + for (std::size_t p : sorted) { + check_position(p); + _selection.erase(_selection.begin() + static_cast(p)); + } + } + + /// Replace the current view with exactly the columns at @p positions, in the + /// given order (projection + reorder in one). Pure index manipulation; + /// allocates nothing. Throws @c std::out_of_range if any position is outside + /// the current view, or @c std::invalid_argument on a duplicate position + /// (the no-alloc materialization moves each column out at most once). + void select_columns(std::span positions) + { + std::vector next; + next.reserve(positions.size()); + for (std::size_t p : positions) { + check_position(p); + next.push_back(_selection[p]); + } + std::vector sorted(next); + std::sort(sorted.begin(), sorted.end()); + if (std::adjacent_find(sorted.begin(), sorted.end()) != sorted.end()) { + throw std::invalid_argument("owning_table_view: duplicate column position in select_columns"); + } + _selection = std::move(next); + } + + /// Realize the currently-exposed columns into an owned @c cudf::table. Moves + /// buffers out of the owner when it is @ref no_alloc_materializable and the + /// base view is column-wise identical to the owner's current view, otherwise + /// copies the selected view (allocating). + [[nodiscard]] std::unique_ptr materialize(rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) + { + return _model->materialize(_selection, stream, mr); + } + + private: + struct owner_concept { + virtual ~owner_concept() = default; + + [[nodiscard]] virtual cudf::table_view full_view() const = 0; + + [[nodiscard]] virtual std::unique_ptr materialize( + std::span selection, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) = 0; + }; + + template + struct owner_model final : owner_concept { + owner_model(Owner owner, cudf::table_view base_view) + : _owner(std::move(owner)), _base_view(std::move(base_view)) + { + } + + [[nodiscard]] cudf::table_view full_view() const override { return _base_view; } + + [[nodiscard]] std::unique_ptr materialize( + std::span selection, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) override + { + if constexpr (no_alloc_materializable) { + // Identity guard: only move columns out when the base view names + // exactly the owner's current columns — otherwise the base_view-relative + // selection indices would pull the wrong columns out of the owner's + // released vector (silent corruption). On mismatch, fall back to copy. + if (shallow_identical(_base_view, (*_owner).view())) { + // Move the selected columns out of the owner; no device buffer copy. + // Precondition: the selection contains no duplicate original indices + // (guaranteed by the swap/drop operations, which only permute/subset an + // identity selection). Moving the same column twice would be a bug, so + // assert each slot is still occupied before moving it out. + auto columns = (*_owner).release(); + std::vector> out; + out.reserve(selection.size()); + for (auto idx : selection) { + auto& column = columns[static_cast(idx)]; + assert(column != nullptr && "owning_table_view: duplicate column index in selection"); + out.push_back(std::move(column)); + } + return std::make_unique(std::move(out)); + } + } + return materialize_by_copy(selection, stream, mr); + } + + /// Generic path: copy the selected view into a fresh table (allocates from @p mr). + [[nodiscard]] std::unique_ptr materialize_by_copy( + std::span selection, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) const + { + std::vector sel(selection.begin(), selection.end()); + return std::make_unique(_base_view.select(sel), stream, mr); + } + + Owner _owner; + cudf::table_view _base_view; + }; + + /// Validate a caller-supplied column position against the current view. + void check_position(std::size_t position) const + { + if (position >= _selection.size()) { + throw std::out_of_range("owning_table_view: column position " + std::to_string(position) + + " is out of range (n_columns=" + std::to_string(_selection.size()) + + ")"); + } + } + + static std::vector identity_selection(cudf::size_type n) + { + std::vector v(static_cast(n)); + std::iota(v.begin(), v.end(), 0); + return v; + } + + std::unique_ptr _model; + std::vector _selection; +}; + +} // namespace detail + +//===----------------------------------------------------------------------===// +// owning_table_view +//===----------------------------------------------------------------------===// +/** + * @brief A handle that exposes a @c cudf::table_view while owning the data + * behind it, regardless of whether that data is a fully-materialized + * @c cudf::table or a view into some other type-erased owner. + * + * The handle is in one of three states: an owned @c cudf::table, a + * @ref detail::owner_backed_view (type-erased owner + column selection), or + * empty. + * + * Only @ref materialize and @ref release may allocate device memory (the + * view -> table transition). Every other mutating operation — + * @ref reorder_columns and @ref drop_columns — is a pure view manipulation: + * if the handle currently owns a @c cudf::table it is first demoted into a + * @ref detail::owner_backed_view (a @c unique_ptr move, not a copy), after + * which the column selection is permuted/subset in place. Because of this they + * are @c const and operate on a @c mutable state member. + */ +class owning_table_view { + public: + /// Construct empty (no valid state). + owning_table_view() = default; + + /// Take ownership of an already-materialized table. A null @p table is + /// normalized to the empty state. + explicit owning_table_view(std::unique_ptr table) + { + if (table) { _state = std::move(table); } + } + + /// Take ownership of a table by move. + explicit owning_table_view(cudf::table&& table) + : _state(std::make_unique(std::move(table))) + { + } + + /** + * @brief Construct from a type-erased @p owner that keeps the data alive and + * an explicit @p base_view naming its columns. + * + * If @p Owner is @ref detail::no_alloc_materializable, materialization will + * move column buffers out of the owner (provided @p base_view names exactly + * the owner's columns); otherwise it copies @p base_view. + */ + template + requires(!std::same_as, owning_table_view> && + !std::same_as, std::unique_ptr>) + owning_table_view(Owner&& owner, cudf::table_view base_view) + : _state(std::make_unique(std::forward(owner), base_view)) + { + } + + /// The currently-exposed view. Empty view when the handle has no valid state. + [[nodiscard]] cudf::table_view view() const; + + /// Number of columns in the currently-exposed view. + [[nodiscard]] std::size_t n_columns() const; + + /// Number of rows in the data. Independent of the exposed column selection; + /// 0 when the handle has no valid state. + [[nodiscard]] cudf::size_type num_rows() const; + + /// The column at the given current position. Throws @c std::out_of_range if + /// @p index is outside the currently-exposed view. + [[nodiscard]] cudf::column_view column(std::size_t index) const; + + /// The data types of the currently-exposed columns, in order. + [[nodiscard]] std::vector column_types() const; + + /// Swap columns at the given current positions (applied in order). Never + /// allocates; demotes an owned table to a view first if necessary. + void reorder_columns(std::span> swaps) const; + + /// Drop columns at the given current positions. Never allocates; demotes an + /// owned table to a view first if necessary. + void drop_columns(std::span positions) const; + + /// Keep exactly the columns at the given current positions, in the given + /// order (projection + reorder). Never allocates; demotes an owned table to a + /// view first if necessary. Throws on out-of-range or duplicate positions. + void select_columns(std::span positions) const; + + /// Realize a view state into an owned table. No-op if already materialized or + /// empty. May allocate (copying path) depending on the underlying owner. + void materialize(rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + + /// Materialize if needed and surrender the owned table, leaving the handle + /// empty. Returns nullptr if the handle had no valid state. + [[nodiscard]] std::unique_ptr release( + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + + /// Release all held data, leaving the handle empty. + void drop(); + + /// True when the handle holds an owned @c cudf::table (materialized state). + [[nodiscard]] bool is_materialized() const; + + /// Allocation size in bytes of the owned table, or 0 when not materialized. + /// The size of a view state is only known once materialized. + [[nodiscard]] std::size_t alloc_size() const; + + /// True when the handle holds a valid state (a table or a view). + explicit operator bool() const { return !std::holds_alternative(_state); } + + private: + /// Demote an owned-table state into an equivalent view state. No-op for the + /// other states. Allocates nothing (moves the table's unique_ptr). + void ensure_view() const; + + mutable std::variant, + std::unique_ptr, + std::monostate> + _state{std::monostate{}}; +}; + +} // namespace cucascade diff --git a/src/cudf/CMakeLists.txt b/src/cudf/CMakeLists.txt index 3302dd9..c05bbaf 100644 --- a/src/cudf/CMakeLists.txt +++ b/src/cudf/CMakeLists.txt @@ -18,6 +18,7 @@ target_sources( cucascade_cudf_objects PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/gpu_data_representation.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/owning_table_view.cpp ${CMAKE_CURRENT_SOURCE_DIR}/host_data_representation.cpp ${CMAKE_CURRENT_SOURCE_DIR}/bandwidth_profiler.cpp ${CMAKE_CURRENT_SOURCE_DIR}/representation_converter_builtins.cpp) diff --git a/src/cudf/gpu_data_representation.cpp b/src/cudf/gpu_data_representation.cpp index d56985a..3be5c2d 100644 --- a/src/cudf/gpu_data_representation.cpp +++ b/src/cudf/gpu_data_representation.cpp @@ -28,7 +28,9 @@ namespace cucascade { gpu_table_representation::gpu_table_representation(std::unique_ptr table, cucascade::memory::memory_space& memory_space, rmm::cuda_stream_view writer_stream) - : idata_representation(memory_space), _table(std::move(table)) + : idata_representation(memory_space), + _alloc_size(table ? table->alloc_size() : 0), + _table(std::move(table)) { // STREAM-LINEAGE: record the writer event in the constructor body so every // representation is born with a recorded event. Skipping when the caller @@ -49,55 +51,58 @@ gpu_table_representation::~gpu_table_representation() } } -std::size_t gpu_table_representation::get_size_in_bytes() const -{ - if (std::holds_alternative>(_table)) { - return std::get>(_table)->alloc_size(); - } else if (std::holds_alternative(_table)) { - return std::get(_table).alloc_size; - } - return 0; -} +std::size_t gpu_table_representation::get_size_in_bytes() const { return _alloc_size; } std::size_t gpu_table_representation::get_uncompressed_data_size_in_bytes() const { return get_size_in_bytes(); } -cudf::table_view gpu_table_representation::get_table_view() const +cudf::table_view gpu_table_representation::get_table_view() const { return _table.view(); } + +std::unique_ptr gpu_table_representation::release_table(rmm::cuda_stream_view stream) { - if (std::holds_alternative>(_table)) { - return std::get>(_table)->view(); - } else { - return std::get(_table).view; - } + // STREAM-LINEAGE: order the (potentially materializing) read after the + // recorded writer event before touching the underlying buffers on `stream`. + if (_writer_event != nullptr) { cucascade::cuda::cuda_event_view{_writer_event}.wait(stream); } + auto table = _table.release(stream, get_memory_space().get_default_allocator()); + _alloc_size = 0; + return table; } -std::unique_ptr gpu_table_representation::release_table( - [[maybe_unused]] rmm::cuda_stream_view stream) +void gpu_table_representation::materialize_table(rmm::cuda_stream_view stream) { - if (std::holds_alternative(_table)) { - _table = std::make_unique(std::get(_table).view, stream); - } - return std::move(std::get>(_table)); + if (!_table || _table.is_materialized()) { return; } + // STREAM-LINEAGE: wait-then-record — `stream` must observe the producing + // writes before the materializing read (move or copy), and the writer event + // is re-recorded afterwards so future readers order against the + // materialization. When no writer event was recorded (legacy paths), the + // caller must pass a stream already ordered after the producing writes. + if (_writer_event != nullptr) { cucascade::cuda::cuda_event_view{_writer_event}.wait(stream); } + _table.materialize(stream, get_memory_space().get_default_allocator()); + _alloc_size = _table.alloc_size(); + record_writer_event(stream); } void gpu_table_representation::rebind_stream(rmm::cuda_stream_view stream) { - // Only the owned-table alternative can be rebound: the owning_table_view alternative - // references memory owned by an external (type-erased) owner, which manages its own + // Only the owned-table state can be rebound: a view state references memory + // owned by an external (type-erased) owner, which manages its own // deallocation stream. - if (!std::holds_alternative>(_table)) { return; } - auto& table = std::get>(_table); - if (!table || table->num_columns() == 0) { return; } - - // cudf::table move-assignment is deleted, so release the columns, rebind each, and rebuild - // the table in place. No device memory is copied and no kernels are launched. - auto columns = table->release(); - for (auto& col : columns) { - col = cudf::rebind_stream(std::move(*col), stream); + if (!_table.is_materialized()) { return; } + // Pure move — the materialized state surrenders its table without touching + // device memory (stream/mr are unused on this path). + auto table = _table.release(stream, get_memory_space().get_default_allocator()); + if (table->num_columns() > 0) { + // cudf::table move-assignment is deleted, so release the columns, rebind each, and rebuild + // the table in place. No device memory is copied and no kernels are launched. + auto columns = table->release(); + for (auto& col : columns) { + col = cudf::rebind_stream(std::move(*col), stream); + } + table = std::make_unique(std::move(columns)); } - table = std::make_unique(std::move(columns)); + _table = owning_table_view{std::move(table)}; } std::unique_ptr gpu_table_representation::clone(rmm::cuda_stream_view stream) diff --git a/src/cudf/owning_table_view.cpp b/src/cudf/owning_table_view.cpp new file mode 100644 index 0000000..0fb3551 --- /dev/null +++ b/src/cudf/owning_table_view.cpp @@ -0,0 +1,143 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include +#include +#include +#include + +namespace cucascade { + +void owning_table_view::ensure_view() const +{ + if (auto* table = std::get_if>(&_state)) { + // Compute the view before moving the table into the type-erased owner. + auto base = (*table)->view(); + _state = std::make_unique(std::move(*table), base); + } +} + +cudf::table_view owning_table_view::view() const +{ + if (auto* table = std::get_if>(&_state)) { return (*table)->view(); } + if (auto* view = std::get_if>(&_state)) { + return (*view)->view(); + } + return cudf::table_view{}; +} + +std::size_t owning_table_view::n_columns() const +{ + if (auto* table = std::get_if>(&_state)) { + return static_cast((*table)->num_columns()); + } + if (auto* view = std::get_if>(&_state)) { + return (*view)->n_columns(); + } + return 0; +} + +cudf::size_type owning_table_view::num_rows() const +{ + if (auto* table = std::get_if>(&_state)) { + return (*table)->num_rows(); + } + if (auto* view = std::get_if>(&_state)) { + return (*view)->num_rows(); + } + return 0; +} + +cudf::column_view owning_table_view::column(std::size_t index) const +{ + auto current = view(); + if (index >= static_cast(current.num_columns())) { + throw std::out_of_range("owning_table_view::column: index " + std::to_string(index) + + " is out of range (n_columns=" + std::to_string(current.num_columns()) + + ")"); + } + return current.column(static_cast(index)); +} + +std::vector owning_table_view::column_types() const +{ + auto current = view(); + std::vector types; + types.reserve(static_cast(current.num_columns())); + for (cudf::size_type i = 0; i < current.num_columns(); ++i) { + types.push_back(current.column(i).type()); + } + return types; +} + +void owning_table_view::reorder_columns( + std::span> swaps) const +{ + if (std::holds_alternative(_state)) { return; } + ensure_view(); + std::get>(_state)->reorder_columns(swaps); +} + +void owning_table_view::drop_columns(std::span positions) const +{ + if (std::holds_alternative(_state)) { return; } + ensure_view(); + std::get>(_state)->drop_columns(positions); +} + +void owning_table_view::select_columns(std::span positions) const +{ + if (std::holds_alternative(_state)) { return; } + ensure_view(); + std::get>(_state)->select_columns(positions); +} + +void owning_table_view::materialize(rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) +{ + if (auto* view = std::get_if>(&_state)) { + _state = (*view)->materialize(stream, mr); + } +} + +std::unique_ptr owning_table_view::release(rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + materialize(stream, mr); + std::unique_ptr out; + if (auto* table = std::get_if>(&_state)) { out = std::move(*table); } + _state = std::monostate{}; + return out; +} + +void owning_table_view::drop() { _state = std::monostate{}; } + +bool owning_table_view::is_materialized() const +{ + return std::holds_alternative>(_state); +} + +std::size_t owning_table_view::alloc_size() const +{ + if (auto* table = std::get_if>(&_state)) { + return (*table)->alloc_size(); + } + return 0; +} + +} // namespace cucascade diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 083428c..992f77a 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -67,6 +67,7 @@ if(NOT CUCASCADE_TOPOLOGY_ONLY AND CUCASCADE_BUILD_CUDF) data/test_data_representation.cpp data/test_disk_host_converters.cpp data/test_gpu_disk_converters.cpp + data/test_owning_table_view.cpp data/test_representation_converter.cpp # Main test runner unittest.cpp) diff --git a/test/data/test_data_representation.cpp b/test/data/test_data_representation.cpp index db4f1e5..5f7f84c 100644 --- a/test/data/test_data_representation.cpp +++ b/test/data/test_data_representation.cpp @@ -2358,3 +2358,140 @@ TEST_CASE("host_data_representation::slice round-trip preserves selected columns std::out_of_range); } } + +// ============================================================================= +// gpu_table_representation materialize_table / release_table +// ============================================================================= + +namespace { + +/// Move-only owner satisfying detail::no_alloc_materializable — exercises the +/// zero-copy move path of the templated gpu_table_representation ctor. +struct table_owner { + std::unique_ptr t; + cudf::table& operator*() const { return *t; } +}; + +/// Device data pointer of each column of a view (buffer identity). +std::vector column_heads(cudf::table_view view) +{ + std::vector out(static_cast(view.num_columns())); + for (cudf::size_type i = 0; i < view.num_columns(); ++i) { + out[static_cast(i)] = view.column(i).head(); + } + return out; +} + +} // namespace + +TEST_CASE("gpu_table_representation materialize_table copies generic-owner views in place", + "[gpu_data_representation][materialize_table]") +{ + auto gpu_space = make_mock_memory_space(memory::Tier::GPU, 0); + rmm::cuda_stream stream; + + auto table = create_simple_cudf_table(64, 2, gpu_space->get_default_allocator(), stream.view()); + auto shared = std::make_shared(std::move(table)); + auto keep = shared; // keeps the source data alive for comparison + const std::size_t estimate = keep->alloc_size() + 4096; // deliberately off + + gpu_table_representation repr( + keep->view(), std::move(shared), estimate, *gpu_space, stream.view()); + REQUIRE(repr.get_size_in_bytes() == estimate); + const auto source_heads = column_heads(keep->view()); + + repr.materialize_table(stream.view()); + stream.synchronize(); + + // shared_ptr owners are copy-constructible, so materialization copies: fresh + // buffers, size updated from the estimate to the actual allocation, writer + // event recorded, contents preserved. + REQUIRE(repr.get_writer_event() != nullptr); + REQUIRE(repr.get_size_in_bytes() != estimate); + REQUIRE(repr.get_size_in_bytes() > 0); + const auto materialized_heads = column_heads(repr.get_table_view()); + for (std::size_t i = 0; i < source_heads.size(); ++i) { + REQUIRE(materialized_heads[i] != source_heads[i]); + } + cucascade::test::expect_cudf_tables_equal_on_stream( + keep->view(), repr.get_table_view(), stream.view()); + + // A second call is a no-op on the now-owned table. + const std::size_t materialized_size = repr.get_size_in_bytes(); + repr.materialize_table(stream.view()); + REQUIRE(repr.get_size_in_bytes() == materialized_size); + REQUIRE(column_heads(repr.get_table_view()) == materialized_heads); + + // Size accounting is consistent with the released table. + auto released = repr.release_table(stream.view()); + REQUIRE(released != nullptr); + REQUIRE(released->alloc_size() == materialized_size); +} + +TEST_CASE("gpu_table_representation materialize_table is a no-op for owned tables", + "[gpu_data_representation][materialize_table]") +{ + auto gpu_space = make_mock_memory_space(memory::Tier::GPU, 0); + rmm::cuda_stream stream; + + auto table = create_simple_cudf_table(32, 2, gpu_space->get_default_allocator(), stream.view()); + auto owned = std::make_unique(std::move(table)); + const auto heads_before = column_heads(owned->view()); + + gpu_table_representation repr(std::move(owned), *gpu_space, stream.view()); + const std::size_t size_before = repr.get_size_in_bytes(); + + repr.materialize_table(stream.view()); + + REQUIRE(repr.get_size_in_bytes() == size_before); + REQUIRE(column_heads(repr.get_table_view()) == heads_before); +} + +TEST_CASE("gpu_table_representation materialize_table moves buffers out of move-only owners", + "[gpu_data_representation][materialize_table]") +{ + auto gpu_space = make_mock_memory_space(memory::Tier::GPU, 0); + rmm::cuda_stream stream; + + auto table = create_simple_cudf_table(48, 2, gpu_space->get_default_allocator(), stream.view()); + auto owned = std::make_unique(std::move(table)); + const auto view = owned->view(); + const auto alloc_size = owned->alloc_size(); + const auto heads_before = column_heads(view); + + gpu_table_representation repr( + view, table_owner{std::move(owned)}, alloc_size, *gpu_space, stream.view()); + + repr.materialize_table(stream.view()); + + // Exclusively-owned, full-view owner: zero-copy — buffer identity preserved. + REQUIRE(column_heads(repr.get_table_view()) == heads_before); + REQUIRE(repr.get_size_in_bytes() == alloc_size); + + auto released = repr.release_table(stream.view()); + REQUIRE(released != nullptr); + REQUIRE(column_heads(released->view()) == heads_before); +} + +TEST_CASE("gpu_table_representation release_table leaves a defined empty state", + "[gpu_data_representation][materialize_table]") +{ + auto gpu_space = make_mock_memory_space(memory::Tier::GPU, 0); + rmm::cuda_stream stream; + + auto table = create_simple_cudf_table(16, 2, gpu_space->get_default_allocator(), stream.view()); + gpu_table_representation repr( + std::make_unique(std::move(table)), *gpu_space, stream.view()); + REQUIRE(repr.get_size_in_bytes() > 0); + + auto released = repr.release_table(stream.view()); + REQUIRE(released != nullptr); + + // After release: zero size, empty view, further releases return nullptr, + // and materialize_table is a harmless no-op. + REQUIRE(repr.get_size_in_bytes() == 0); + REQUIRE(repr.get_table_view().num_columns() == 0); + REQUIRE(repr.release_table(stream.view()) == nullptr); + repr.materialize_table(stream.view()); + REQUIRE(repr.get_size_in_bytes() == 0); +} diff --git a/test/data/test_owning_table_view.cpp b/test/data/test_owning_table_view.cpp new file mode 100644 index 0000000..02f8289 --- /dev/null +++ b/test/data/test_owning_table_view.cpp @@ -0,0 +1,374 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include +#include +#include +#include + +using namespace cucascade; + +namespace { + +rmm::cuda_stream_view test_stream() { return cudf::get_default_stream(); } +rmm::device_async_resource_ref test_mr() { return cudf::get_current_device_resource_ref(); } + +// Build an INT32 column whose single value tags the column for identification. +std::unique_ptr tagged_column(std::int32_t tag) +{ + auto col = cudf::make_numeric_column( + cudf::data_type{cudf::type_id::INT32}, 1, cudf::mask_state::UNALLOCATED, test_stream()); + cudaMemcpy(col->mutable_view().data(), &tag, sizeof(tag), cudaMemcpyHostToDevice); + return col; +} + +// Build a table with one INT32 column per tag. +std::unique_ptr tagged_table(std::vector const& tags) +{ + std::vector> cols; + cols.reserve(tags.size()); + for (auto t : tags) { + cols.push_back(tagged_column(t)); + } + return std::make_unique(std::move(cols)); +} + +// Read back the single tag value of each column of a view. +std::vector read_tags(cudf::table_view view) +{ + std::vector out(static_cast(view.num_columns())); + for (cudf::size_type i = 0; i < view.num_columns(); ++i) { + cudaMemcpy(&out[static_cast(i)], + view.column(i).data(), + sizeof(std::int32_t), + cudaMemcpyDeviceToHost); + } + return out; +} + +// The device data pointers of each column of a view (column identity). +std::vector data_ptrs(cudf::table_view view) +{ + std::vector out(static_cast(view.num_columns())); + for (cudf::size_type i = 0; i < view.num_columns(); ++i) { + out[static_cast(i)] = view.column(i).head(); + } + return out; +} + +// A keep-alive owner that does NOT satisfy no_alloc_materializable (no +// dereferenceable release()), forcing the copying materialization path. +struct table_keepalive { + std::unique_ptr table; +}; + +// A move-only owner that satisfies no_alloc_materializable but whose exposed +// view() is NOT the base_view it was paired with — the identity guard must +// reject the move and materialize the base view by copy instead. +struct decoy_owner { + std::unique_ptr real; ///< keeps the base_view's data alive + std::unique_ptr decoy; ///< what (*this) dereferences to + cudf::table& operator*() const { return *decoy; } +}; + +// The concept gates the zero-copy move path: exclusively-owned, +// column-surrendering owners only. +static_assert(detail::no_alloc_materializable>); +static_assert(detail::no_alloc_materializable); +static_assert(!detail::no_alloc_materializable>, + "shared owners must copy — moving would gut co-owners"); +static_assert(!detail::no_alloc_materializable); + +} // namespace + +TEST_CASE("owning_table_view empty state", "[owning_table_view]") +{ + owning_table_view handle; + REQUIRE_FALSE(static_cast(handle)); + REQUIRE(handle.n_columns() == 0); + REQUIRE(handle.view().num_columns() == 0); + + // Mutating an empty handle is a no-op and leaves it empty. + std::array drop{0}; + handle.drop_columns(drop); + REQUIRE_FALSE(static_cast(handle)); + + // release() on an empty handle yields nullptr. + REQUIRE(handle.release(test_stream(), test_mr()) == nullptr); +} + +TEST_CASE("owning_table_view from owned table exposes the table", "[owning_table_view]") +{ + owning_table_view handle{tagged_table({10, 20, 30})}; + REQUIRE(static_cast(handle)); + REQUIRE(handle.n_columns() == 3); + REQUIRE(read_tags(handle.view()) == std::vector{10, 20, 30}); +} + +TEST_CASE("owning_table_view num_rows / column / column_types accessors", "[owning_table_view]") +{ + // 4 columns, 1 row each. + owning_table_view handle{tagged_table({10, 20, 30, 40})}; + + REQUIRE(handle.num_rows() == 1); + REQUIRE(handle.column_types().size() == 4); + for (auto const& t : handle.column_types()) { + REQUIRE(t.id() == cudf::type_id::INT32); + } + + // column() returns the column at the current position; row count is + // unaffected by dropping columns. + std::array drop{0}; + handle.drop_columns(drop); // -> 20, 30, 40 + REQUIRE(handle.num_rows() == 1); + REQUIRE(handle.column_types().size() == 3); + + std::int32_t tag = 0; + cudaMemcpy(&tag, handle.column(0).data(), sizeof(tag), cudaMemcpyDeviceToHost); + REQUIRE(tag == 20); + + REQUIRE_THROWS_AS(handle.column(3), std::out_of_range); + + // Empty handle: zero rows, no columns. + owning_table_view empty; + REQUIRE(empty.num_rows() == 0); + REQUIRE(empty.column_types().empty()); + REQUIRE_THROWS_AS(empty.column(0), std::out_of_range); +} + +TEST_CASE("owning_table_view drop + materialize moves buffers (no copy)", "[owning_table_view]") +{ + auto table = tagged_table({10, 20, 30}); + // Capture the original device pointers before handing the table over. + auto original_ptrs = data_ptrs(table->view()); + + owning_table_view handle{std::move(table)}; + + // Drop the middle column — pure view manipulation, no allocation. + std::array drop{1}; + handle.drop_columns(drop); + REQUIRE(handle.n_columns() == 2); + REQUIRE(read_tags(handle.view()) == std::vector{10, 30}); + + // Materialize back to an owned table: surviving columns keep their buffers. + auto result = handle.release(test_stream(), test_mr()); + REQUIRE(result != nullptr); + REQUIRE(read_tags(result->view()) == std::vector{10, 30}); + + auto result_ptrs = data_ptrs(result->view()); + REQUIRE(result_ptrs[0] == original_ptrs[0]); // column 10 buffer preserved + REQUIRE(result_ptrs[1] == original_ptrs[2]); // column 30 buffer preserved + + // Handle is now empty. + REQUIRE_FALSE(static_cast(handle)); + REQUIRE(handle.release(test_stream(), test_mr()) == nullptr); +} + +TEST_CASE("owning_table_view reorder swaps columns", "[owning_table_view]") +{ + owning_table_view handle{tagged_table({10, 20, 30})}; + + // Swap positions 0 and 2. + std::array, 1> swaps{{{0, 2}}}; + handle.reorder_columns(swaps); + REQUIRE(read_tags(handle.view()) == std::vector{30, 20, 10}); + + auto result = handle.release(test_stream(), test_mr()); + REQUIRE(read_tags(result->view()) == std::vector{30, 20, 10}); +} + +TEST_CASE("owning_table_view reorder then drop compose on current positions", "[owning_table_view]") +{ + owning_table_view handle{tagged_table({10, 20, 30, 40})}; + + std::array, 1> swaps{{{0, 3}}}; + handle.reorder_columns(swaps); // -> 40, 20, 30, 10 + std::array drop{1, 2}; + handle.drop_columns(drop); // drop current positions 1 and 2 -> 40, 10 + REQUIRE(handle.n_columns() == 2); + REQUIRE(read_tags(handle.view()) == std::vector{40, 10}); +} + +TEST_CASE("owning_table_view select_columns projects and reorders (no copy)", "[owning_table_view]") +{ + auto table = tagged_table({10, 20, 30, 40}); + auto original_ptrs = data_ptrs(table->view()); + + owning_table_view handle{std::move(table)}; + + // Project down to columns 3 and 0, in that order. + std::array keep{3, 0}; + handle.select_columns(keep); + REQUIRE(handle.n_columns() == 2); + REQUIRE(read_tags(handle.view()) == std::vector{40, 10}); + + // select_columns composes on current positions: now keep only position 1 (10). + std::array keep2{1}; + handle.select_columns(keep2); + REQUIRE(read_tags(handle.view()) == std::vector{10}); + + // Zero-alloc: the surviving buffer is the original column 10 buffer. + auto result = handle.release(test_stream(), test_mr()); + REQUIRE(data_ptrs(result->view())[0] == original_ptrs[0]); +} + +TEST_CASE("owning_table_view select_columns rejects bad positions", "[owning_table_view]") +{ + owning_table_view handle{tagged_table({10, 20, 30})}; + + std::array out_of_range{0, 3}; + REQUIRE_THROWS_AS(handle.select_columns(out_of_range), std::out_of_range); + + std::array duplicate{1, 1}; + REQUIRE_THROWS_AS(handle.select_columns(duplicate), std::invalid_argument); + + // Unchanged after rejected operations. + REQUIRE(handle.n_columns() == 3); + REQUIRE(read_tags(handle.view()) == std::vector{10, 20, 30}); +} + +TEST_CASE("owning_table_view rejects out-of-range positions", "[owning_table_view]") +{ + owning_table_view handle{tagged_table({10, 20, 30})}; + + std::array bad_drop{3}; + REQUIRE_THROWS_AS(handle.drop_columns(bad_drop), std::out_of_range); + + std::array, 1> bad_swap{{{0, 3}}}; + REQUIRE_THROWS_AS(handle.reorder_columns(bad_swap), std::out_of_range); + + // The handle is unchanged after the rejected operations. + REQUIRE(handle.n_columns() == 3); + REQUIRE(read_tags(handle.view()) == std::vector{10, 20, 30}); +} + +TEST_CASE("owning_table_view generic owner materializes by copy", "[owning_table_view]") +{ + auto table = tagged_table({10, 20, 30}); + auto original_ptrs = data_ptrs(table->view()); + auto base_view = table->view(); + + // Type-erase a non-no_alloc owner; materialization must copy. + owning_table_view handle{table_keepalive{std::move(table)}, base_view}; + REQUIRE(handle.n_columns() == 3); + + std::array drop{1}; + handle.drop_columns(drop); + REQUIRE(read_tags(handle.view()) == std::vector{10, 30}); + + auto result = handle.release(test_stream(), test_mr()); + REQUIRE(read_tags(result->view()) == std::vector{10, 30}); + + // Copying path: the materialized buffers differ from the originals. + auto result_ptrs = data_ptrs(result->view()); + REQUIRE(result_ptrs[0] != original_ptrs[0]); + REQUIRE(result_ptrs[1] != original_ptrs[2]); +} + +TEST_CASE("owning_table_view is_materialized and alloc_size track state", "[owning_table_view]") +{ + owning_table_view empty; + REQUIRE_FALSE(empty.is_materialized()); + REQUIRE(empty.alloc_size() == 0); + + owning_table_view handle{tagged_table({10, 20, 30})}; + REQUIRE(handle.is_materialized()); + REQUIRE(handle.alloc_size() > 0); + + // Column ops demote the owned table to a view state: size no longer known. + std::array drop{0}; + handle.drop_columns(drop); + REQUIRE(static_cast(handle)); + REQUIRE_FALSE(handle.is_materialized()); + REQUIRE(handle.alloc_size() == 0); + + handle.materialize(test_stream(), test_mr()); + REQUIRE(handle.is_materialized()); + REQUIRE(handle.alloc_size() > 0); + + auto out = handle.release(test_stream(), test_mr()); + REQUIRE(out != nullptr); + REQUIRE_FALSE(handle.is_materialized()); + REQUIRE(handle.alloc_size() == 0); +} + +TEST_CASE("owning_table_view null owned table yields an empty handle", "[owning_table_view]") +{ + owning_table_view handle{std::unique_ptr{}}; + REQUIRE_FALSE(static_cast(handle)); + REQUIRE(handle.n_columns() == 0); + REQUIRE(handle.num_rows() == 0); + REQUIRE_FALSE(handle.is_materialized()); + REQUIRE(handle.alloc_size() == 0); + REQUIRE(handle.release(test_stream(), test_mr()) == nullptr); +} + +TEST_CASE("owning_table_view shared owner materializes by copy (co-owner kept intact)", + "[owning_table_view]") +{ + auto table = tagged_table({10, 20}); + auto original_ptrs = data_ptrs(table->view()); + + auto shared = std::shared_ptr(std::move(table)); + auto keep = shared; // co-owner that must not be gutted + owning_table_view handle{std::move(shared), keep->view()}; + + auto result = handle.release(test_stream(), test_mr()); + REQUIRE(read_tags(result->view()) == std::vector{10, 20}); + + // Copying path: fresh buffers, and the co-owner still holds its columns. + auto result_ptrs = data_ptrs(result->view()); + REQUIRE(result_ptrs[0] != original_ptrs[0]); + REQUIRE(result_ptrs[1] != original_ptrs[1]); + REQUIRE(keep->num_columns() == 2); + REQUIRE(read_tags(keep->view()) == std::vector{10, 20}); +} + +TEST_CASE("owning_table_view identity guard falls back to copy on mismatched base view", + "[owning_table_view]") +{ + auto real = tagged_table({10, 20}); + auto decoy = tagged_table({77, 88}); + auto base_view = real->view(); + auto real_ptrs = data_ptrs(base_view); + + // decoy_owner satisfies the move concept, but its view() is the decoy table, + // not base_view — moving would surrender the WRONG columns. The guard must + // detect the mismatch and copy the base view instead. + owning_table_view handle{decoy_owner{std::move(real), std::move(decoy)}, base_view}; + + auto result = handle.release(test_stream(), test_mr()); + REQUIRE(result != nullptr); + REQUIRE(read_tags(result->view()) == std::vector{10, 20}); // base data, not decoy + + auto result_ptrs = data_ptrs(result->view()); + REQUIRE(result_ptrs[0] != real_ptrs[0]); + REQUIRE(result_ptrs[1] != real_ptrs[1]); +}