diff --git a/include/core/image.hpp b/include/core/image.hpp new file mode 100644 index 00000000..ef7707a8 --- /dev/null +++ b/include/core/image.hpp @@ -0,0 +1,177 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#pragma once + +#include +#include +#include +#include +#include + +#include "core/detail/allocators/i_allocator.hpp" +#include "core/image_buffer.hpp" +#include "core/image_data.hpp" +#include "core/image_format.hpp" +#include "core/util_enums.h" +#include "operator_types.h" + +namespace roccv { + +class ImageStorage; + +/** + * @brief Cleanup callback signature for ImageWrapData. Invoked when the last + * Image handle referencing the wrapped buffer is destroyed. Receives the + * ImageData snapshot that was originally wrapped, so callbacks can free + * multi-plane buffers or dispatch on format. + */ +using ImageDataCleanupFunc = std::function; + +/** + * @brief Per-image allocation spec describing what to allocate for a single + * variable-sized image. Mirrors NVCVImageRequirements: size, format, per-plane + * row strides, and base-address alignment. Used as the input to Image's + * allocating constructors and as the output of CalcRequirements; also + * preserved on the Image itself as the source of truth from which exportData() + * rebuilds an ImageData snapshot on demand. + * + * Per-plane row strides are populated only for planes 0..numPlanes(format)-1; + * remaining slots are unused. Today's interleaved-only ImageFormat means only + * planeRowStride[0] is populated in practice. + */ +struct ImageRequirements { + Size2D size; // Width and height in pixels. + ImageFormat format; // Pixel format (dtype + channel count + swizzle). + int64_t planeRowStride[ROCCV_MAX_IMAGE_PLANES]; // Per-plane row stride in bytes. + int32_t alignBytes; // Required base-address alignment, in bytes. +}; + +/** + * @brief A single variable-sized image with device-resident pixel data. + * + * Image is the per-element type held by ImageBatchVarShape. It is a handle + * over a refcounted ImageStorage: copying an Image bumps the refcount and + * leaves both handles pointing at the same underlying buffer. The buffer is + * freed when the last handle is destroyed (for owning Images) or when the + * cleanup callback fires (for ImageWrapData with a callback). + * + * Storage shape: Image holds the buffer pointer (via ImageStorage) plus the + * "ingredients" describing it (size, format, device, per-plane row strides). + * It does NOT hold a precomputed ImageData snapshot — exportData() rebuilds + * one on demand from the ingredients. This keeps a single source of truth for + * the buffer pointer and aligns with how ImageBatchVarShape produces its + * own snapshots. + */ +class Image { + public: + using Requirements = ImageRequirements; + + /** + * @brief Compute the requirements (row stride, etc.) for an image of the + * given dimensions and format. + */ + static Requirements CalcRequirements(Size2D size, ImageFormat format); + + /** + * @brief Allocate a new device buffer for an image of the given dimensions + * and format using the global default allocator. + */ + explicit Image(Size2D size, ImageFormat format, eDeviceType device = eDeviceType::GPU); + + /** + * @brief Allocate a new device buffer using a caller-supplied allocator. + */ + explicit Image(Size2D size, ImageFormat format, const IAllocator& alloc, eDeviceType device = eDeviceType::GPU); + + /** + * @brief Allocate a new device buffer from precomputed requirements. + */ + explicit Image(const Requirements& reqs, eDeviceType device = eDeviceType::GPU); + explicit Image(const Requirements& reqs, const IAllocator& alloc, eDeviceType device = eDeviceType::GPU); + + Image(const Image&) = default; // refcount bump + Image(Image&&) noexcept = default; + Image& operator=(const Image&) = default; // refcount bump + Image& operator=(Image&&) noexcept = default; + ~Image() = default; + + Size2D size() const noexcept { return m_size; } + ImageFormat format() const noexcept { return m_format; } + eDeviceType device() const noexcept { return m_device; } + + /** + * @brief Build and return an ImageData snapshot describing this image. + * + * Returned by value (not by reference) — Image stores ingredients, not a + * precomputed snapshot, so each call constructs a fresh ImageData. The + * snapshot's plane descriptors point into this Image's buffer; it remains + * valid as long as any handle to this storage is alive. + */ + ImageData exportData() const; + + /** + * @brief Build a snapshot and down-cast it to a specific subclass. Throws + * std::bad_cast if the underlying buffer kind doesn't match Derived. + */ + template + Derived exportData() const; + + private: + Image(const Requirements& reqs, eDeviceType device, std::shared_ptr storage); + + friend Image ImageWrapData(const ImageData& data, ImageDataCleanupFunc cleanup); + + std::shared_ptr m_data; + Size2D m_size; + ImageFormat m_format; + eDeviceType m_device; + std::array m_planeRowStride; +}; + +template +Derived Image::exportData() const { + ImageData data = exportData(); + auto derived = data.cast(); + if (!derived.has_value()) { + throw std::bad_cast(); + } + return derived.value(); +} + +/** + * @brief Wrap an externally-owned buffer as an Image without allocating. + * + * View-only by default: the wrapped buffer is NOT freed when the returned + * Image (and any copies) go out of scope. The caller is responsible for + * keeping the underlying memory alive for as long as any handle survives. + * + * Pass a non-null cleanup callback to opt into ownership transfer; the + * callback runs exactly once, when the last handle is destroyed. + * + * @param[in] data Pre-existing image data (pointer, layout, device). + * @param[in] cleanup Optional callback to free the buffer on last destruction. + * @return An Image referencing the wrapped buffer. + */ +extern Image ImageWrapData(const ImageData& data, ImageDataCleanupFunc cleanup = nullptr); + +} // namespace roccv diff --git a/include/core/image_batch_buffer.hpp b/include/core/image_batch_buffer.hpp new file mode 100644 index 00000000..06f0c4b3 --- /dev/null +++ b/include/core/image_batch_buffer.hpp @@ -0,0 +1,96 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#pragma once + +#include + +#include "core/image_buffer.hpp" +#include "core/image_format.hpp" + +namespace roccv { + +/** + * @brief Pitch-linear descriptor table for a variable-shape image batch. + * + * Each entry of `imageList` is a full per-image strided buffer descriptor — + * reusing `ImageBufferStrided` keeps the per-image shape (multi-plane-capable, + * one base pointer per plane, per-plane row stride) identical to what a single + * `Image` carries today. + * + * Pointer residency: + * - `imageList` is the descriptor table read by GPU kernels. For a GPU-resident + * batch this points into device memory; for a hypothetical CPU-resident + * batch it would point into host memory. The producing batch class owns the + * allocation and decides residency. + * - `formatList` mirrors `imageList`'s residency and holds one ImageFormat per + * image (so kernels can branch on per-image format without dereferencing the + * descriptor table). + * - `hostFormatList` is always host-resident. It exists so host-side validation + * code can read per-image formats without paying a D->H copy. For a + * CPU-resident batch this MAY alias `formatList`; for a GPU-resident batch + * it is a separate host mirror kept in sync by the producer. + * + * `uniqueFormat` is the common ImageFormat across all images, or FMT_NONE if + * formats are heterogeneous or the batch is empty. Cached to fast-path the + * homogeneous case. + * + * `maxWidth` / `maxHeight` are the bounding box across all images. Used by + * operators to size launch grids. Both are 0 when the batch is empty. + * + * The struct is intentionally trivially copyable so it can ride inside + * `ImageBatchBuffer` without an allocation, mirroring `ImageBufferStrided`'s + * relationship to `ImageBuffer`. + */ +struct ImageBatchVarShapeBufferStrided { + /** Common format across all images in the batch, or a default-constructed + * ImageFormat if formats are heterogeneous or the batch is empty. */ + ImageFormat uniqueFormat; + + /** Bounding box across all images, in pixels. Both 0 when empty. */ + int32_t maxWidth; + int32_t maxHeight; + + /** Per-image format array, length == numImages. Residency matches + * `imageList` (device for GPU batches, host for CPU batches). */ + ImageFormat* formatList; + + /** Host-resident mirror of `formatList`. May alias `formatList` for + * CPU-resident batches. Length == numImages. */ + const ImageFormat* hostFormatList; + + /** Per-image descriptor table, length == numImages. The kernel-facing + * pointer; residency determines which device the batch lives on. */ + ImageBufferStrided* imageList; +}; + +/** + * @brief An image-batch buffer. Currently only the variable-shape strided + * variant is supported. Shaped as a tagged-union-style aggregate so additional + * batch buffer kinds can be added later (e.g. tensor-backed batches) without + * changing the public type. + */ +struct ImageBatchBuffer { + ImageBatchVarShapeBufferStrided varShapeStrided; +}; + +} // namespace roccv diff --git a/include/core/image_batch_data.hpp b/include/core/image_batch_data.hpp new file mode 100644 index 00000000..c50de010 --- /dev/null +++ b/include/core/image_batch_data.hpp @@ -0,0 +1,244 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#pragma once + +#include + +#include +#include + +#include "core/image_batch_buffer.hpp" +#include "core/image_format.hpp" +#include "core/util_enums.h" +#include "operator_types.h" + +namespace roccv { + +/** + * @brief Discriminator for the kind of buffer an ImageBatchData carries. Used + * by IsCompatibleKind() / cast<>() to perform safe runtime down-casting through + * the ImageBatchData hierarchy. + * + * The hierarchy currently exposes only one concrete buffer kind + * (variable-shape, strided, GPU-resident); the enum is shaped to grow into + * additional kinds (e.g. tensor-backed batches, host-resident varshape) without + * breaking the existing buffer kind values. + */ +enum class ImageBatchBufferType { + IMAGE_BATCH_BUFFER_NONE, // Default/invalid buffer type. + IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP, // GPU-accessible varshape descriptor table. + IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST, // Host-accessible varshape descriptor table. +}; + +/** + * @brief Holds the underlying image-batch data alongside metadata + * (numImages, buffer kind). Non-strided batch data is not supported for use + * right now; use ImageBatchVarShapeDataStrided to access strided varshape data + * instead. + * + * ImageBatchData is the interchange type for a batch of variable-sized images. + * It does not own any of the underlying buffers (the descriptor table, the + * format arrays, or the per-image pixel buffers) — it is a metadata snapshot, + * valid only as long as the producing batch outlives it. + * + * Lazy-sync note: for a GPU-resident batch the producer (ImageBatchVarShape) + * is responsible for ensuring the device-side descriptor table is up to date + * with any pushBack/popBack edits before handing out an ImageBatchData. The + * snapshot itself carries no synchronization state. + */ +class ImageBatchData { + public: + ImageBatchData() = delete; + virtual ~ImageBatchData() = default; + + /** + * @brief Returns the number of images currently in the batch. + */ + virtual int32_t numImages() const; + + /** + * @brief Returns the device the descriptor table (and per-image pixel + * buffers) reside on. + */ + virtual eDeviceType device() const; + + /** + * @brief Attempts to down-cast this ImageBatchData to a more specific + * subclass. Returns the casted value if the underlying buffer kind matches + * what Derived expects, or std::nullopt otherwise. + * + * @tparam Derived The target subclass to cast to. + */ + template + std::optional cast() const { + static_assert(std::is_base_of::value, + "Cannot cast ImageBatchData to an unrelated type."); + static_assert(sizeof(Derived) == sizeof(ImageBatchData), + "Derived type must not add any additional data members."); + + if (!Derived::IsCompatibleKind(m_bufferType)) { + return std::nullopt; + } + + return std::make_optional(m_numImages, m_buffer); + } + + static bool IsCompatibleKind(ImageBatchBufferType bufferType); + + protected: + ImageBatchData(int32_t numImages, const ImageBatchBuffer& buffer); + + int32_t m_numImages; + eDeviceType m_deviceType; + ImageBatchBufferType m_bufferType; + ImageBatchBuffer m_buffer; +}; + +/** + * @brief Image-batch data backed by a variable-shape descriptor table. Adds + * typed accessors for the per-image format arrays and the bounding box across + * the batch. Sub-classed by ImageBatchVarShapeDataStrided to discriminate + * pitch-linear storage; further sub-classed by ImageBatchVarShapeDataStridedHip + * to tag device residency. + */ +class ImageBatchVarShapeData : public ImageBatchData { + public: + using Buffer = ImageBatchVarShapeBufferStrided; + + ImageBatchVarShapeData(int32_t numImages, const ImageBatchBuffer& buffer); + + static bool IsCompatibleKind(ImageBatchBufferType bufferType); + + /** + * @brief Bounding box across all images in the batch, in pixels. Both + * dimensions are 0 when the batch is empty. Used by operators to size + * launch grids without iterating the descriptor table. + */ + Size2D maxSize() const; + + /** + * @brief Returns the common ImageFormat across all images, or FMT_NONE if + * formats are heterogeneous or the batch is empty. + */ + ImageFormat uniqueFormat() const; + + /** + * @brief Per-image format array. Residency matches the descriptor table + * (device for GPU batches). Length == numImages(). + * + * Prefer hostFormatList() for host-side validation paths to avoid a D->H + * copy. + */ + const ImageFormat* formatList() const; + + /** + * @brief Host-resident mirror of formatList(). Always safe to dereference + * from host code. Length == numImages(). + */ + const ImageFormat* hostFormatList() const; +}; + +/** + * @brief Variable-shape image-batch data backed by a pitch-linear descriptor + * table. Adds the per-image descriptor accessor on top of + * ImageBatchVarShapeData. + */ +class ImageBatchVarShapeDataStrided : public ImageBatchVarShapeData { + public: + using Buffer = ImageBatchVarShapeBufferStrided; + + ImageBatchVarShapeDataStrided(int32_t numImages, const ImageBatchBuffer& buffer); + + static bool IsCompatibleKind(ImageBatchBufferType bufferType); + + /** + * @brief Per-image descriptor table. Length == numImages(). Residency + * matches the enclosing data type — for ImageBatchVarShapeDataStridedHip + * this is a device pointer; kernels read it directly. + * + * Each entry is a full ImageBufferStrided so the per-image shape + * (multi-plane-capable, per-plane stride and base pointer) matches what a + * single Image carries. + */ + const ImageBufferStrided* imageList() const; +}; + +/** + * @brief GPU-accessible variable-shape image-batch data. + */ +class ImageBatchVarShapeDataStridedHip : public ImageBatchVarShapeDataStrided { + public: + using Buffer = ImageBatchVarShapeBufferStrided; + + ImageBatchVarShapeDataStridedHip(int32_t numImages, const ImageBatchBuffer& buffer); + + /** + * @brief Constructs GPU-accessible varshape image-batch data from the + * concrete strided buffer directly. + * + * @param[in] numImages Number of images currently in the batch. + * @param[in] buffer Descriptor table + per-image format arrays. The + * descriptor table and `formatList` must point to GPU + * memory; `hostFormatList` to host memory. + */ + ImageBatchVarShapeDataStridedHip(int32_t numImages, const Buffer& buffer); + + static bool IsCompatibleKind(ImageBatchBufferType bufferType); +}; + +/** + * @brief Host-accessible variable-shape image-batch data. + * + * The host-resident counterpart to ImageBatchVarShapeDataStridedHip. The + * descriptor table, `formatList`, and `hostFormatList` all point to host + * memory; `formatList` and `hostFormatList` MAY alias the same allocation + * since no D->H sync is required. + * + * The lazy host->device descriptor sync that the GPU producer needs is not + * applicable here — host-only varshape batches can edit the descriptor table + * in place and hand it straight to host kernels. The matching producer-side + * design (whether host batches are a separate type, a runtime-tagged variant + * of ImageBatchVarShape, or skipped entirely in favor of CPU-side per-image + * loops) is still open. + */ +class ImageBatchVarShapeDataStridedHost : public ImageBatchVarShapeDataStrided { + public: + using Buffer = ImageBatchVarShapeBufferStrided; + + ImageBatchVarShapeDataStridedHost(int32_t numImages, const ImageBatchBuffer& buffer); + + /** + * @brief Constructs host-accessible varshape image-batch data from the + * concrete strided buffer directly. + * + * @param[in] numImages Number of images currently in the batch. + * @param[in] buffer Descriptor table + per-image format arrays. All + * pointers must reference host memory; `formatList` + * and `hostFormatList` may alias. + */ + ImageBatchVarShapeDataStridedHost(int32_t numImages, const Buffer& buffer); + + static bool IsCompatibleKind(ImageBatchBufferType bufferType); +}; + +} // namespace roccv diff --git a/include/core/image_batch_var_shape.hpp b/include/core/image_batch_var_shape.hpp new file mode 100644 index 00000000..ba4fd147 --- /dev/null +++ b/include/core/image_batch_var_shape.hpp @@ -0,0 +1,206 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#pragma once + +#include +#include + +#include +#include +#include + +#include "core/detail/allocators/i_allocator.hpp" +#include "core/image.hpp" +#include "core/image_batch_data.hpp" +#include "core/image_format.hpp" +#include "exception.hpp" +#include "operator_types.h" + +namespace roccv { + +/** + * @brief Producer-side container for a batch of variable-sized images that + * share a single GPU-resident descriptor table. + * + * Holds up to `capacity()` Image handles and maintains a parallel descriptor + * table that operators can dispatch over without iterating Image-by-Image. + * Capacity is fixed at construction; pushBack/popBack move within it. + * + * The host descriptor mirrors are pinned so the H2D copy in exportData() is a + * true DMA (no runtime bounce buffer) and so the snapshot can expose the same + * pinned pointer as both `formatList`'s host shadow and `hostFormatList`. + * + * Sync model: pushBack/popBack mutate the host mirrors only; the device + * descriptor table is brought up to date lazily inside exportData(stream), + * which copies just the dirty suffix `[dirtyStart, numImages)`. A hipEvent + * (`m_postFence`) guards the host buffers — if a previous exportData's H2D + * is still in flight, pushBack hipEventSynchronize's on the CPU before + * mutating, so the snapshot a consumer is reading never tears. + * + * GPU-only in v1. CPU-resident images are rejected on push. + */ +class ImageBatchVarShape { + public: + using const_iterator = std::vector::const_iterator; + + /** + * @brief Construct an empty batch with `capacity` slots, using the global + * default allocator. + */ + explicit ImageBatchVarShape(int32_t capacity); + + /** + * @brief Construct an empty batch with `capacity` slots, using the supplied + * allocator. The allocator must outlive the batch. + */ + explicit ImageBatchVarShape(int32_t capacity, const IAllocator &alloc); + + ~ImageBatchVarShape(); + + ImageBatchVarShape(const ImageBatchVarShape &) = delete; + ImageBatchVarShape &operator=(const ImageBatchVarShape &) = delete; + ImageBatchVarShape(ImageBatchVarShape &&) noexcept; + ImageBatchVarShape &operator=(ImageBatchVarShape &&) = delete; + + int32_t capacity() const noexcept { return m_capacity; } + int32_t numImages() const noexcept { return static_cast(m_images.size()); } + + /** + * @brief Append an image to the batch. Throws if capacity would be + * exceeded, the image is CPU-resident, or the image has more than one + * plane (rocCV is single-plane today). + */ + void pushBack(const Image &img); + + /** + * @brief Append a range of images. Strong exception guarantee — if any + * image fails validation, the batch is rolled back to its pre-call state + * and the exception is rethrown. + */ + template + void pushBack(It begin, It end); + + /** + * @brief Remove the trailing `count` images. Throws if `count` exceeds + * numImages(). + */ + void popBack(int32_t count = 1); + + /** + * @brief Drop all images. Buffers are kept; the batch is reusable. + */ + void clear(); + + const Image &operator[](int32_t i) const { return m_images[i]; } + + const_iterator begin() const noexcept { return m_images.cbegin(); } + const_iterator end() const noexcept { return m_images.cend(); } + + /** + * @brief Bounding box across all images, in pixels. Returns Size2D{0, 0} + * for an empty batch. + */ + Size2D maxSize() const; + + /** + * @brief The common ImageFormat across all images, or FMT_NONE if formats + * are heterogeneous or the batch is empty. popBack invalidates the cache + * so the next call rescans and may return an exact format again. + */ + ImageFormat uniqueFormat() const; + + /** + * @brief Build (and return by value) a GPU-resident snapshot of the batch. + * + * Synchronizes the dirty suffix of the host mirrors to the device + * descriptor table on the supplied stream before returning. The returned + * snapshot's `imageList` and `formatList` are device pointers safe for + * kernels enqueued on the same stream; `hostFormatList` aliases the pinned + * host format mirror and is safe to read from host code. The snapshot is + * a metadata view valid as long as this batch outlives it. + */ + ImageBatchVarShapeDataStridedHip exportData(hipStream_t stream); + + /** + * @brief Build a snapshot and down-cast it to a specific subclass. Throws + * std::bad_cast if the underlying buffer kind doesn't match Derived. + */ + template + Derived exportData(hipStream_t stream); + + private: + void doSyncDirtySuffix(hipStream_t stream); + void doUpdateCache() const; + + int32_t m_capacity; + int32_t m_dirtyStartingFromIndex = 0; + bool m_fencePending = false; + + const IAllocator &m_allocator; + std::vector m_images; + + ImageBufferStrided *m_devImagesBuffer = nullptr; + ImageFormat *m_devFormatsBuffer = nullptr; + ImageBufferStrided *m_hostImagesBuffer = nullptr; + ImageFormat *m_hostFormatsBuffer = nullptr; + + hipEvent_t m_postFence = nullptr; + + mutable std::optional m_cacheMaxSize; + mutable std::optional m_cacheUniqueFormat; +}; + +template +void ImageBatchVarShape::pushBack(It begin, It end) { + const int32_t incoming = static_cast(std::distance(begin, end)); + if (incoming + numImages() > m_capacity) { + throw Exception("ImageBatchVarShape::pushBack range would exceed capacity", eStatusType::OUT_OF_BOUNDS); + } + + const int32_t oldNumImages = numImages(); + const auto oldMaxSize = m_cacheMaxSize; + const auto oldUniqueFormat = m_cacheUniqueFormat; + + try { + for (auto it = begin; it != end; ++it) { + pushBack(*it); + } + } catch (...) { + m_images.erase(m_images.begin() + oldNumImages, m_images.end()); + m_cacheMaxSize = oldMaxSize; + m_cacheUniqueFormat = oldUniqueFormat; + throw; + } +} + +template +Derived ImageBatchVarShape::exportData(hipStream_t stream) { + ImageBatchVarShapeDataStridedHip data = exportData(stream); + auto derived = data.cast(); + if (!derived.has_value()) { + throw std::bad_cast(); + } + return derived.value(); +} + +} // namespace roccv diff --git a/include/core/image_buffer.hpp b/include/core/image_buffer.hpp new file mode 100644 index 00000000..aea93c2d --- /dev/null +++ b/include/core/image_buffer.hpp @@ -0,0 +1,81 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#pragma once + +#include + +/** Maximum number of data planes an image can have. */ +#define ROCCV_MAX_IMAGE_PLANES (6) + +namespace roccv { + +/** + * @brief Describes a single pitch-linear image plane. + * + * For interleaved-channel formats there is exactly one plane covering the whole + * image. For planar formats (e.g. NV12, YUV420) each channel/plane carries its + * own width, height, and row stride and lives in its own buffer. + */ +struct ImagePlaneStrided { + /** Width of this plane in pixels. Must be >= 1. */ + int32_t width; + + /** Height of this plane in pixels. Must be >= 1. */ + int32_t height; + + /** Distance in bytes between the start of consecutive rows. Must be at + * least `(width * bits-per-pixel + 7) / 8`. */ + int64_t rowStride; + + /** Pointer to the first byte of plane data. Validity (device vs host) is + * determined by the enclosing data type. */ + void* basePtr; +}; + +/** + * @brief A pitch-linear image buffer: one or more `ImagePlaneStrided` entries. + * + * Only the first `numPlanes` entries carry valid data; the remainder of the + * fixed-size `planes` array is unused. Capping the array size keeps the buffer + * trivially copyable so it can ride inside `ImageBuffer` without an + * allocation. + */ +struct ImageBufferStrided { + /** Number of valid planes. Must be >= 1. */ + int32_t numPlanes; + + /** Per-plane descriptors. Only the first `numPlanes` are valid. */ + ImagePlaneStrided planes[ROCCV_MAX_IMAGE_PLANES]; +}; + +/** + * @brief An image buffer. Currently only the strided variant is supported. + * Mirrors the role `TensorBuffer` plays for tensors and is intentionally + * shaped as a tagged-union-style aggregate so additional buffer kinds can be + * added later (e.g. HIP textures) without changing the public type. + */ +struct ImageBuffer { + ImageBufferStrided strided; +}; + +} // namespace roccv diff --git a/include/core/image_data.hpp b/include/core/image_data.hpp new file mode 100644 index 00000000..cf45e71c --- /dev/null +++ b/include/core/image_data.hpp @@ -0,0 +1,178 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#pragma once + +#include + +#include +#include + +#include "core/image_buffer.hpp" +#include "core/image_format.hpp" +#include "core/util_enums.h" +#include "operator_types.h" + +namespace roccv { + +/** + * @brief Discriminator for the kind of buffer an ImageData carries. Used by + * IsCompatibleKind() / cast<>() to perform safe runtime down-casting through + * the ImageData hierarchy. + */ +enum class ImageBufferType { + IMAGE_BUFFER_NONE, // Default/invalid buffer type. Used when no buffer type is specified. + IMAGE_BUFFER_STRIDED_HIP, // GPU-accessible buffer with strided access. + IMAGE_BUFFER_STRIDED_HOST // Host-accessible buffer with strided access. +}; + +/** + * @brief Holds the underlying image data alongside metadata (format, buffer + * kind). Non-strided image data is not supported for use right now; use + * ImageDataStrided to access strided image data instead. + * + * ImageData is the interchange type for a single variable-sized image. It + * does not own the underlying pixel buffer — it is a metadata snapshot, valid + * only as long as the producing buffer outlives it. + */ +class ImageData { + public: + ImageData() = delete; + virtual ~ImageData() = default; + + /** + * @brief Returns the pixel format of the image. + */ + virtual const ImageFormat &format() const; + + /** + * @brief Returns the device the image data resides on. + */ + virtual eDeviceType device() const; + + /** + * @brief Attempts to down-cast this ImageData to a more specific subclass. + * Returns the casted value if the underlying buffer kind matches what + * Derived expects, or std::nullopt otherwise. + * + * @tparam Derived The target subclass to cast to. + */ + template + std::optional cast() const { + static_assert(std::is_base_of::value, "Cannot cast ImageData to an unrelated type."); + static_assert(sizeof(Derived) == sizeof(ImageData), "Derived type must not add any additional data members."); + + if (!Derived::IsCompatibleKind(m_bufferType)) { + return std::nullopt; + } + + return std::make_optional(m_format, m_buffer); + } + + static bool IsCompatibleKind(ImageBufferType bufferType); + + protected: + ImageData(const ImageFormat &format, const ImageBuffer &buffer); + + ImageFormat m_format; + eDeviceType m_deviceType; + ImageBufferType m_bufferType; + ImageBuffer m_buffer; +}; + +/** + * @brief Image data backed by one or more pitch-linear planes. Adds typed + * accessors for plane descriptors on top of the base ImageData. Sub-classed + * by ImageDataStridedHip and ImageDataStridedHost to discriminate device vs + * host residency. + */ +class ImageDataStrided : public ImageData { + public: + using Buffer = ImageBufferStrided; + + ImageDataStrided(const ImageFormat &format, const ImageBuffer &buffer); + + static bool IsCompatibleKind(ImageBufferType bufferType); + + /** + * @brief Returns the logical image dimensions, taken from plane 0. For + * planar formats, individual planes may have smaller dimensions (e.g. + * chroma sub-sampling); use plane(p) to inspect each plane directly. + */ + Size2D size() const; + + /** + * @brief Returns the number of valid planes in the buffer. + */ + int32_t numPlanes() const; + + /** + * @brief Returns the descriptor for the requested plane. + * + * @param[in] p The plane index. Must satisfy `0 <= p < numPlanes()`. + */ + const ImagePlaneStrided &plane(int32_t p) const; +}; + +/** + * @brief GPU-accessible strided image data. + */ +class ImageDataStridedHip : public ImageDataStrided { + public: + using Buffer = ImageBufferStrided; + + ImageDataStridedHip(const ImageFormat &format, const ImageBuffer &buffer); + + /** + * @brief Constructs GPU-accessible strided image data from a strided + * image buffer directly. + * + * @param[in] format The pixel format. + * @param[in] buffer A strided image buffer with planes allocated on the GPU. + */ + ImageDataStridedHip(const ImageFormat &format, const Buffer &buffer); + + static bool IsCompatibleKind(ImageBufferType bufferType); +}; + +/** + * @brief Host-accessible strided image data. + */ +class ImageDataStridedHost : public ImageDataStrided { + public: + using Buffer = ImageBufferStrided; + + ImageDataStridedHost(const ImageFormat &format, const ImageBuffer &buffer); + + /** + * @brief Constructs host-accessible strided image data from a strided + * image buffer directly. + * + * @param[in] format The pixel format. + * @param[in] buffer A strided image buffer with planes allocated on the host. + */ + ImageDataStridedHost(const ImageFormat &format, const Buffer &buffer); + + static bool IsCompatibleKind(ImageBufferType bufferType); +}; + +} // namespace roccv diff --git a/include/core/image_format.hpp b/include/core/image_format.hpp index 7dd891f3..ddb1d100 100644 --- a/include/core/image_format.hpp +++ b/include/core/image_format.hpp @@ -40,7 +40,10 @@ enum class eSwizzle { */ class ImageFormat { public: - explicit ImageFormat() {} + /** + * @brief Default-constructs to FMT_NONE. + */ + constexpr ImageFormat() : m_dtype(eDataType::DATA_TYPE_U8), m_numChannels(0), m_swizzle(eSwizzle::XYZW) {} explicit constexpr ImageFormat(eDataType dtype, int32_t numChannels, eSwizzle swizzle = eSwizzle::XYZW) : m_dtype(dtype), m_numChannels(numChannels), m_swizzle(swizzle) {} @@ -48,12 +51,20 @@ class ImageFormat { int32_t channels() const noexcept; eSwizzle swizzle() const noexcept; + constexpr bool operator==(const ImageFormat& other) const noexcept { + return m_dtype == other.m_dtype && m_numChannels == other.m_numChannels && m_swizzle == other.m_swizzle; + } + constexpr bool operator!=(const ImageFormat& other) const noexcept { return !(*this == other); } + private: eDataType m_dtype; int32_t m_numChannels; eSwizzle m_swizzle; }; +// Undefined format. Used to represent an uninitialized or invalid format. +constexpr ImageFormat FMT_NONE{eDataType::DATA_TYPE_U8, 0, eSwizzle::XYZW}; + // Single plane with one 8-bit unsigned integer channel. constexpr ImageFormat FMT_U8(eDataType::DATA_TYPE_U8, 1, eSwizzle::XYZW); diff --git a/include/core/image_storage.hpp b/include/core/image_storage.hpp new file mode 100644 index 00000000..70984742 --- /dev/null +++ b/include/core/image_storage.hpp @@ -0,0 +1,53 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#pragma once + +namespace roccv { + +/** + * @brief Holds the raw data pointer for a single Image and serves as the + * refcount target shared between Image handles. + * + * ImageStorage carries no lifecycle logic of its own: freeing the underlying + * buffer is the responsibility of the shared_ptr deleter + * installed at the Image construction site. The allocating Image ctor + * captures the allocator + device into its deleter; ImageWrapData captures + * the user's cleanup callback (or installs none for the view-only case). + * + * As a result, ImageStorage is held only by shared_ptr — never by value, never + * copied. Move/copy are deleted to enforce that. + */ +class ImageStorage { + public: + explicit ImageStorage(void* data) : m_data(data) {} + + ImageStorage(const ImageStorage&) = delete; + ImageStorage& operator=(const ImageStorage&) = delete; + + void* data() const noexcept { return m_data; } + + private: + void* m_data; +}; + +} // namespace roccv diff --git a/src/core/image.cpp b/src/core/image.cpp new file mode 100644 index 00000000..d6077dcb --- /dev/null +++ b/src/core/image.cpp @@ -0,0 +1,186 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "core/image.hpp" + +#include +#include + +#include "core/data_type.hpp" +#include "core/detail/context.hpp" +#include "core/exception.hpp" +#include "core/image_storage.hpp" + +namespace roccv { + +namespace { + +// Allocates a buffer through `alloc` for the requested device and wraps it +// in an ImageStorage whose shared_ptr deleter frees through the same allocator. +// The allocator reference is captured by reference; callers must ensure it +// outlives every Image (and any handle copied from it) it creates. +std::shared_ptr makeStorage(const ImageRequirements& reqs, const IAllocator& alloc, eDeviceType device) { + const size_t bytes = static_cast(reqs.planeRowStride[0]) * reqs.size.h; + + void* buf = nullptr; + switch (device) { + case eDeviceType::GPU: + buf = alloc.allocHipMem(bytes); + break; + case eDeviceType::CPU: + buf = alloc.allocHostMem(bytes); + break; + } + + return std::shared_ptr(new ImageStorage(buf), [&alloc, device](ImageStorage* s) { + switch (device) { + case eDeviceType::GPU: + alloc.freeHipMem(s->data()); + break; + case eDeviceType::CPU: + alloc.freeHostMem(s->data()); + break; + } + delete s; + }); +} + +} // namespace + +// ----------------------------------------------------------------------------- +// CalcRequirements +// ----------------------------------------------------------------------------- + +Image::Requirements Image::CalcRequirements(Size2D size, ImageFormat format) { + if (size.w < 1 || size.h < 1) { + throw Exception("Image dimensions must be >= 1.", eStatusType::INVALID_VALUE); + } + + const int64_t bytesPerPixel = static_cast(DataType(format.dtype()).size()) * format.channels(); + + // Guard signed-overflow in the rowStride = bytesPerPixel * width product + // (UB on overflow). Realistic image sizes don't approach INT64_MAX, but + // pathological callers shouldn't silently propagate garbage into strides. + int64_t rowStride = 0; + if (__builtin_mul_overflow(bytesPerPixel, static_cast(size.w), &rowStride)) { + throw Exception("Image row stride overflows int64.", eStatusType::INVALID_VALUE); + } + + // TODO: derive a sensible default base/row alignment from device attributes. + return ImageRequirements{ + .size = size, + .format = format, + .planeRowStride = {rowStride}, + .alignBytes = 0, + }; +} + +// ----------------------------------------------------------------------------- +// Constructors +// ----------------------------------------------------------------------------- + +Image::Image(Size2D size, ImageFormat format, eDeviceType device) + : Image(size, format, GlobalContext().getDefaultAllocator(), device) {} + +Image::Image(Size2D size, ImageFormat format, const IAllocator& alloc, eDeviceType device) + : Image(CalcRequirements(size, format), alloc, device) {} + +Image::Image(const Requirements& reqs, eDeviceType device) + : Image(reqs, GlobalContext().getDefaultAllocator(), device) {} + +Image::Image(const Requirements& reqs, const IAllocator& alloc, eDeviceType device) + : Image(reqs, device, makeStorage(reqs, alloc, device)) {} + +Image::Image(const Requirements& reqs, eDeviceType device, std::shared_ptr storage) + : m_data(std::move(storage)), m_size(reqs.size), m_format(reqs.format), m_device(device), m_planeRowStride{} { + std::copy(std::begin(reqs.planeRowStride), std::end(reqs.planeRowStride), m_planeRowStride.begin()); +} + +// ----------------------------------------------------------------------------- +// exportData +// ----------------------------------------------------------------------------- + +ImageData Image::exportData() const { + // TODO: derive numPlanes from m_format when planar formats land. Today's + // ImageFormat is interleaved-only, so plane 0 covers the whole image and + // its dimensions match m_size verbatim. + ImageBufferStrided strided{}; + strided.numPlanes = 1; + strided.planes[0].width = m_size.w; + strided.planes[0].height = m_size.h; + strided.planes[0].rowStride = m_planeRowStride[0]; + strided.planes[0].basePtr = m_data->data(); + + switch (m_device) { + case eDeviceType::GPU: + return ImageDataStridedHip(m_format, strided); + case eDeviceType::CPU: + return ImageDataStridedHost(m_format, strided); + } + + throw Exception("Unsupported device type in Image::exportData.", eStatusType::INVALID_VALUE); +} + +// ----------------------------------------------------------------------------- +// ImageWrapData +// ----------------------------------------------------------------------------- + +Image ImageWrapData(const ImageData& data, ImageDataCleanupFunc cleanup) { + auto strided = data.cast(); + if (!strided.has_value()) { + throw Exception("ImageWrapData requires strided image data.", eStatusType::INVALID_VALUE); + } + + // Single-plane assumption: storage tracks plane(0) and Requirements only + // populates planeRowStride[0]. Multi-plane wraps will need to copy each + // plane's stride and either store per-plane base pointers or derive them + // from a single owning allocation. + const ImagePlaneStrided& plane0 = strided->plane(0); + + // Designated initializers to avoid value-initializing ImageFormat through + // its explicit default ctor (which copy-list-init refuses). + Image::Requirements reqs{ + .size = Size2D{plane0.width, plane0.height}, + .format = data.format(), + .planeRowStride = {plane0.rowStride}, + .alignBytes = 0, + }; + + // The deleter captures `data` by value so the original snapshot survives + // long enough to be passed to the cleanup callback on last-handle drop. + // Swallow exceptions from `cleanup` — shared_ptr deleters run during + // destruction, and a throw would propagate into std::terminate. + auto storage = + std::shared_ptr(new ImageStorage(plane0.basePtr), [data, cleanup](ImageStorage* s) noexcept { + if (cleanup) { + try { + cleanup(data); + } catch (...) { + } + } + delete s; + }); + + return Image(reqs, data.device(), std::move(storage)); +} + +} // namespace roccv diff --git a/src/core/image_batch_data.cpp b/src/core/image_batch_data.cpp new file mode 100644 index 00000000..a8ce07ba --- /dev/null +++ b/src/core/image_batch_data.cpp @@ -0,0 +1,103 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "core/image_batch_data.hpp" + +#include "core/image_batch_buffer.hpp" +#include "core/image_format.hpp" +#include "core/util_enums.h" + +namespace roccv { + +int32_t ImageBatchData::numImages() const { return m_numImages; } + +eDeviceType ImageBatchData::device() const { return m_deviceType; } + +ImageBatchData::ImageBatchData(int32_t numImages, const ImageBatchBuffer& buffer) + : m_numImages(numImages), + m_deviceType(eDeviceType::GPU), + m_bufferType(ImageBatchBufferType::IMAGE_BATCH_BUFFER_NONE), + m_buffer(buffer) {} + +bool ImageBatchData::IsCompatibleKind(ImageBatchBufferType bufferType) { + return bufferType != ImageBatchBufferType::IMAGE_BATCH_BUFFER_NONE; +} + +ImageBatchVarShapeData::ImageBatchVarShapeData(int32_t numImages, const ImageBatchBuffer& buffer) + : ImageBatchData(numImages, buffer) {} + +bool ImageBatchVarShapeData::IsCompatibleKind(ImageBatchBufferType bufferType) { + return bufferType == ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP || + bufferType == ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST; +} + +Size2D ImageBatchVarShapeData::maxSize() const { + return Size2D{m_buffer.varShapeStrided.maxWidth, m_buffer.varShapeStrided.maxHeight}; +} + +ImageFormat ImageBatchVarShapeData::uniqueFormat() const { return m_buffer.varShapeStrided.uniqueFormat; } + +const ImageFormat* ImageBatchVarShapeData::formatList() const { return m_buffer.varShapeStrided.formatList; } + +const ImageFormat* ImageBatchVarShapeData::hostFormatList() const { return m_buffer.varShapeStrided.hostFormatList; } + +ImageBatchVarShapeDataStrided::ImageBatchVarShapeDataStrided(int32_t numImages, const ImageBatchBuffer& buffer) + : ImageBatchVarShapeData(numImages, buffer) {} + +bool ImageBatchVarShapeDataStrided::IsCompatibleKind(ImageBatchBufferType bufferType) { + return bufferType == ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP || + bufferType == ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST; +} + +const ImageBufferStrided* ImageBatchVarShapeDataStrided::imageList() const { + return m_buffer.varShapeStrided.imageList; +} + +ImageBatchVarShapeDataStridedHip::ImageBatchVarShapeDataStridedHip(int32_t numImages, const ImageBatchBuffer& buffer) + : ImageBatchVarShapeDataStrided(numImages, buffer) { + m_bufferType = ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP; + m_deviceType = eDeviceType::GPU; +} + +ImageBatchVarShapeDataStridedHip::ImageBatchVarShapeDataStridedHip( + int32_t numImages, const ImageBatchVarShapeDataStridedHip::Buffer& buffer) + : ImageBatchVarShapeDataStridedHip(numImages, ImageBatchBuffer{.varShapeStrided = buffer}) {} + +bool ImageBatchVarShapeDataStridedHip::IsCompatibleKind(ImageBatchBufferType bufferType) { + return bufferType == ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP; +} + +ImageBatchVarShapeDataStridedHost::ImageBatchVarShapeDataStridedHost(int32_t numImages, const ImageBatchBuffer& buffer) + : ImageBatchVarShapeDataStrided(numImages, buffer) { + m_bufferType = ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST; + m_deviceType = eDeviceType::CPU; +} + +ImageBatchVarShapeDataStridedHost::ImageBatchVarShapeDataStridedHost( + int32_t numImages, const ImageBatchVarShapeDataStridedHost::Buffer& buffer) + : ImageBatchVarShapeDataStridedHost(numImages, ImageBatchBuffer{.varShapeStrided = buffer}) {} + +bool ImageBatchVarShapeDataStridedHost::IsCompatibleKind(ImageBatchBufferType bufferType) { + return bufferType == ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST; +} + +} // namespace roccv diff --git a/src/core/image_batch_var_shape.cpp b/src/core/image_batch_var_shape.cpp new file mode 100644 index 00000000..510cccac --- /dev/null +++ b/src/core/image_batch_var_shape.cpp @@ -0,0 +1,252 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "core/image_batch_var_shape.hpp" + +#include + +#include "core/detail/context.hpp" +#include "core/exception.hpp" +#include "core/hip_assert.h" +#include "core/image_batch_buffer.hpp" +#include "core/image_buffer.hpp" + +namespace roccv { + +ImageBatchVarShape::ImageBatchVarShape(int32_t capacity) + : ImageBatchVarShape(capacity, GlobalContext().getDefaultAllocator()) {} + +ImageBatchVarShape::ImageBatchVarShape(int32_t capacity, const IAllocator& alloc) + : m_capacity(capacity), m_allocator(alloc) { + if (capacity <= 0) { + throw Exception("ImageBatchVarShape capacity must be positive", eStatusType::INVALID_VALUE); + } + + m_images.reserve(capacity); + + const size_t imagesBytes = sizeof(ImageBufferStrided) * capacity; + const size_t formatsBytes = sizeof(ImageFormat) * capacity; + + try { + m_devImagesBuffer = static_cast(m_allocator.allocHipMem(imagesBytes)); + m_devFormatsBuffer = static_cast(m_allocator.allocHipMem(formatsBytes)); + m_hostImagesBuffer = static_cast(m_allocator.allocHostPinnedMem(imagesBytes)); + m_hostFormatsBuffer = static_cast(m_allocator.allocHostPinnedMem(formatsBytes)); + + HIP_VALIDATE_NO_ERRORS(hipEventCreateWithFlags(&m_postFence, hipEventDisableTiming)); + } catch (...) { + if (m_hostFormatsBuffer != nullptr) m_allocator.freeHostPinnedMem(m_hostFormatsBuffer); + if (m_hostImagesBuffer != nullptr) m_allocator.freeHostPinnedMem(m_hostImagesBuffer); + if (m_devFormatsBuffer != nullptr) m_allocator.freeHipMem(m_devFormatsBuffer); + if (m_devImagesBuffer != nullptr) m_allocator.freeHipMem(m_devImagesBuffer); + throw; + } +} + +ImageBatchVarShape::~ImageBatchVarShape() { + if (m_fencePending && m_postFence != nullptr) { + // Drain any in-flight H2D copy before freeing the host mirrors it + // reads from. (void) — destructors must not throw. + (void)hipEventSynchronize(m_postFence); + } + if (m_postFence != nullptr) { + (void)hipEventDestroy(m_postFence); + } + if (m_hostFormatsBuffer != nullptr) m_allocator.freeHostPinnedMem(m_hostFormatsBuffer); + if (m_hostImagesBuffer != nullptr) m_allocator.freeHostPinnedMem(m_hostImagesBuffer); + if (m_devFormatsBuffer != nullptr) m_allocator.freeHipMem(m_devFormatsBuffer); + if (m_devImagesBuffer != nullptr) m_allocator.freeHipMem(m_devImagesBuffer); +} + +ImageBatchVarShape::ImageBatchVarShape(ImageBatchVarShape&& other) noexcept + : m_capacity(other.m_capacity), + m_dirtyStartingFromIndex(other.m_dirtyStartingFromIndex), + m_fencePending(other.m_fencePending), + m_allocator(other.m_allocator), + m_images(std::move(other.m_images)), + m_devImagesBuffer(other.m_devImagesBuffer), + m_devFormatsBuffer(other.m_devFormatsBuffer), + m_hostImagesBuffer(other.m_hostImagesBuffer), + m_hostFormatsBuffer(other.m_hostFormatsBuffer), + m_postFence(other.m_postFence), + m_cacheMaxSize(other.m_cacheMaxSize), + m_cacheUniqueFormat(other.m_cacheUniqueFormat) { + other.m_capacity = 0; + other.m_dirtyStartingFromIndex = 0; + other.m_fencePending = false; + other.m_devImagesBuffer = nullptr; + other.m_devFormatsBuffer = nullptr; + other.m_hostImagesBuffer = nullptr; + other.m_hostFormatsBuffer = nullptr; + other.m_postFence = nullptr; + other.m_cacheMaxSize.reset(); + other.m_cacheUniqueFormat.reset(); +} + +void ImageBatchVarShape::pushBack(const Image& img) { + const int32_t n = numImages(); + if (n >= m_capacity) { + throw Exception("ImageBatchVarShape::pushBack would exceed capacity", eStatusType::OUT_OF_BOUNDS); + } + if (img.device() != eDeviceType::GPU) { + throw Exception("ImageBatchVarShape only accepts GPU-resident images", eStatusType::INVALID_VALUE); + } + + ImageDataStridedHip data = img.exportData(); + if (data.numPlanes() != 1) { + throw Exception("ImageBatchVarShape only supports single-plane images", eStatusType::INVALID_VALUE); + } + + if (m_fencePending) { + HIP_VALIDATE_NO_ERRORS(hipEventSynchronize(m_postFence)); + m_fencePending = false; + } + + ImageBufferStrided slot{}; + slot.numPlanes = 1; + slot.planes[0] = data.plane(0); + m_hostImagesBuffer[n] = slot; + m_hostFormatsBuffer[n] = img.format(); + + const Size2D imgSize = img.size(); + if (n == 0) { + // Seed from scratch: an empty-batch query may have populated the + // cache with sentinels (FMT_NONE, 0×0); replacing avoids merging the + // first real image into them. + m_cacheMaxSize = imgSize; + m_cacheUniqueFormat = img.format(); + } else { + // popBack invalidates m_cacheMaxSize without rescanning, so make sure + // both halves of the cache are populated before merging in. + doUpdateCache(); + m_cacheMaxSize->w = std::max(m_cacheMaxSize->w, imgSize.w); + m_cacheMaxSize->h = std::max(m_cacheMaxSize->h, imgSize.h); + if (*m_cacheUniqueFormat != img.format()) { + m_cacheUniqueFormat = FMT_NONE; + } + } + + m_images.push_back(img); +} + +void ImageBatchVarShape::popBack(int32_t count) { + if (count < 0) { + throw Exception("ImageBatchVarShape::popBack count must be non-negative", eStatusType::INVALID_VALUE); + } + if (count > numImages()) { + throw Exception("ImageBatchVarShape::popBack count exceeds numImages", eStatusType::OUT_OF_BOUNDS); + } + + m_images.erase(m_images.end() - count, m_images.end()); + m_dirtyStartingFromIndex = std::min(m_dirtyStartingFromIndex, numImages()); + + // maxSize can only shrink on pop; force a rescan on next query. uniqueFormat + // stays — it may now be conservatively FMT_NONE, but never wrong. + m_cacheMaxSize.reset(); + if (numImages() == 0) { + m_cacheUniqueFormat.reset(); + } +} + +void ImageBatchVarShape::clear() { + m_images.clear(); + m_dirtyStartingFromIndex = 0; + m_cacheMaxSize.reset(); + m_cacheUniqueFormat.reset(); +} + +Size2D ImageBatchVarShape::maxSize() const { + doUpdateCache(); + return m_cacheMaxSize.value_or(Size2D{0, 0}); +} + +ImageFormat ImageBatchVarShape::uniqueFormat() const { + doUpdateCache(); + return m_cacheUniqueFormat.value_or(FMT_NONE); +} + +void ImageBatchVarShape::doUpdateCache() const { + if (m_cacheMaxSize.has_value() && m_cacheUniqueFormat.has_value()) { + return; + } + const int32_t n = static_cast(m_images.size()); + if (n == 0) { + m_cacheMaxSize = Size2D{0, 0}; + m_cacheUniqueFormat = FMT_NONE; + return; + } + + Size2D maxSz{0, 0}; + ImageFormat unique = m_hostFormatsBuffer[0]; + bool heterogeneous = false; + for (int32_t i = 0; i < n; ++i) { + const ImagePlaneStrided& p0 = m_hostImagesBuffer[i].planes[0]; + maxSz.w = std::max(maxSz.w, p0.width); + maxSz.h = std::max(maxSz.h, p0.height); + if (!heterogeneous && m_hostFormatsBuffer[i] != unique) { + heterogeneous = true; + } + } + m_cacheMaxSize = maxSz; + m_cacheUniqueFormat = heterogeneous ? FMT_NONE : unique; +} + +void ImageBatchVarShape::doSyncDirtySuffix(hipStream_t stream) { + const int32_t n = numImages(); + if (m_dirtyStartingFromIndex >= n) { + return; + } + const int32_t dirtyCount = n - m_dirtyStartingFromIndex; + + if (m_fencePending) { + HIP_VALIDATE_NO_ERRORS(hipStreamWaitEvent(stream, m_postFence, /*flags=*/0)); + } + + HIP_VALIDATE_NO_ERRORS(hipMemcpyAsync(m_devImagesBuffer + m_dirtyStartingFromIndex, + m_hostImagesBuffer + m_dirtyStartingFromIndex, + sizeof(ImageBufferStrided) * dirtyCount, hipMemcpyHostToDevice, stream)); + HIP_VALIDATE_NO_ERRORS(hipMemcpyAsync(m_devFormatsBuffer + m_dirtyStartingFromIndex, + m_hostFormatsBuffer + m_dirtyStartingFromIndex, + sizeof(ImageFormat) * dirtyCount, hipMemcpyHostToDevice, stream)); + + HIP_VALIDATE_NO_ERRORS(hipEventRecord(m_postFence, stream)); + m_fencePending = true; + m_dirtyStartingFromIndex = n; +} + +ImageBatchVarShapeDataStridedHip ImageBatchVarShape::exportData(hipStream_t stream) { + doSyncDirtySuffix(stream); + doUpdateCache(); + + const Size2D maxSz = m_cacheMaxSize.value(); + ImageBatchVarShapeBufferStrided buffer{}; + buffer.uniqueFormat = m_cacheUniqueFormat.value(); + buffer.maxWidth = maxSz.w; + buffer.maxHeight = maxSz.h; + buffer.formatList = m_devFormatsBuffer; + buffer.hostFormatList = m_hostFormatsBuffer; + buffer.imageList = m_devImagesBuffer; + + return ImageBatchVarShapeDataStridedHip(numImages(), buffer); +} + +} // namespace roccv diff --git a/src/core/image_data.cpp b/src/core/image_data.cpp new file mode 100644 index 00000000..6fb0fc83 --- /dev/null +++ b/src/core/image_data.cpp @@ -0,0 +1,88 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "core/image_data.hpp" + +#include "core/image_buffer.hpp" +#include "core/image_format.hpp" +#include "core/util_enums.h" + +namespace roccv { + +const ImageFormat& ImageData::format() const { return m_format; } + +eDeviceType ImageData::device() const { return m_deviceType; } + +ImageData::ImageData(const ImageFormat& format, const ImageBuffer& buffer) + : m_format(format), + m_deviceType(eDeviceType::GPU), + m_bufferType(ImageBufferType::IMAGE_BUFFER_NONE), + m_buffer(buffer) {} + +bool ImageData::IsCompatibleKind(ImageBufferType bufferType) { + return bufferType != ImageBufferType::IMAGE_BUFFER_NONE; +} + +ImageDataStrided::ImageDataStrided(const ImageFormat& format, const ImageBuffer& buffer) + : ImageData(format, buffer) {} + +bool ImageDataStrided::IsCompatibleKind(ImageBufferType bufferType) { + return bufferType == ImageBufferType::IMAGE_BUFFER_STRIDED_HIP || + bufferType == ImageBufferType::IMAGE_BUFFER_STRIDED_HOST; +} + +Size2D ImageDataStrided::size() const { + const ImagePlaneStrided& p0 = m_buffer.strided.planes[0]; + return Size2D{p0.width, p0.height}; +} + +int32_t ImageDataStrided::numPlanes() const { return m_buffer.strided.numPlanes; } + +const ImagePlaneStrided& ImageDataStrided::plane(int32_t p) const { return m_buffer.strided.planes[p]; } + +ImageDataStridedHip::ImageDataStridedHip(const ImageFormat& format, const ImageBuffer& buffer) + : ImageDataStrided(format, buffer) { + m_bufferType = ImageBufferType::IMAGE_BUFFER_STRIDED_HIP; + m_deviceType = eDeviceType::GPU; +} + +ImageDataStridedHip::ImageDataStridedHip(const ImageFormat& format, const ImageDataStridedHip::Buffer& buffer) + : ImageDataStridedHip(format, ImageBuffer{.strided = buffer}) {} + +bool ImageDataStridedHip::IsCompatibleKind(ImageBufferType bufferType) { + return bufferType == ImageBufferType::IMAGE_BUFFER_STRIDED_HIP; +} + +ImageDataStridedHost::ImageDataStridedHost(const ImageFormat& format, const ImageBuffer& buffer) + : ImageDataStrided(format, buffer) { + m_bufferType = ImageBufferType::IMAGE_BUFFER_STRIDED_HOST; + m_deviceType = eDeviceType::CPU; +} + +ImageDataStridedHost::ImageDataStridedHost(const ImageFormat& format, const ImageDataStridedHost::Buffer& buffer) + : ImageDataStridedHost(format, ImageBuffer{.strided = buffer}) {} + +bool ImageDataStridedHost::IsCompatibleKind(ImageBufferType bufferType) { + return bufferType == ImageBufferType::IMAGE_BUFFER_STRIDED_HOST; +} + +} // namespace roccv diff --git a/tests/roccv/cpp/include/image_test_helpers.hpp b/tests/roccv/cpp/include/image_test_helpers.hpp new file mode 100644 index 00000000..c4613367 --- /dev/null +++ b/tests/roccv/cpp/include/image_test_helpers.hpp @@ -0,0 +1,129 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ +#pragma once + +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace roccv { +namespace tests { + +// Opaque sentinel pointers used by image-layer tests. ImageData / ImageBatchData +// carry pointers but never dereference them — the buffer is a metadata snapshot +// only — so tests use these to verify values flow through without needing real +// allocations. +inline void* const FAKE_PTR_A = reinterpret_cast(0xAAAAAAAAull); +inline void* const FAKE_PTR_B = reinterpret_cast(0xBBBBBBBBull); +inline void* const FAKE_PTR_C = reinterpret_cast(0xCCCCCCCCull); + +/** + * @brief Test allocator that backs every allocation kind with malloc and tallies + * how many times each entry point is invoked. Pure host-backed; no actual GPU + * dependency on the returned pointers — callers that exercise the Hip/pinned + * paths must only inspect metadata, never dereference device memory. + * + * `lastAllocBytes` is updated from every alloc path (hip, host, pinned), so + * callers may assert on the most recent allocation regardless of kind. + */ +class CountingAllocator : public IAllocator { + public: + mutable int hipAllocs = 0; + mutable int hipFrees = 0; + mutable int hostAllocs = 0; + mutable int hostFrees = 0; + mutable int pinnedAllocs = 0; + mutable int pinnedFrees = 0; + mutable size_t lastAllocBytes = 0; + + void* allocHipMem(size_t size) const override { + ++hipAllocs; + lastAllocBytes = size; + return std::malloc(size); + } + void freeHipMem(void* ptr) const noexcept override { + ++hipFrees; + std::free(ptr); + } + + void* allocHostMem(size_t size, int32_t /*alignment*/ = 0) const override { + ++hostAllocs; + lastAllocBytes = size; + return std::malloc(size); + } + void freeHostMem(void* ptr) const noexcept override { + ++hostFrees; + std::free(ptr); + } + + void* allocHostPinnedMem(size_t size) const override { + ++pinnedAllocs; + lastAllocBytes = size; + return std::malloc(size); + } + void freeHostPinnedMem(void* ptr) const noexcept override { + ++pinnedFrees; + std::free(ptr); + } +}; + +// Single-plane packed-row buffer descriptor around `basePtr`. The pointer is +// never dereferenced by the consumers (ImageData / ImageBatchVarShape). +inline ImageBufferStrided MakeSinglePlaneBuffer(int32_t width, int32_t height, int64_t rowStride, void* basePtr) { + ImageBufferStrided buf{}; + buf.numPlanes = 1; + buf.planes[0] = {width, height, rowStride, basePtr}; + return buf; +} + +// Single-plane GPU-resident ImageData snapshot with packed-row stride implied +// by `fmt`. For tests that need an ImageData but won't touch the pixels. +inline ImageDataStridedHip MakeFakeHipData(int32_t width, int32_t height, void* basePtr, ImageFormat fmt = FMT_RGB8) { + const int64_t rowStride = static_cast(width) * fmt.channels() * DataType(fmt.dtype()).size(); + return ImageDataStridedHip(fmt, MakeSinglePlaneBuffer(width, height, rowStride, basePtr)); +} + +// Host counterpart of MakeFakeHipData. +inline ImageDataStridedHost MakeFakeHostData(int32_t width, int32_t height, void* basePtr, ImageFormat fmt = FMT_RGB8) { + const int64_t rowStride = static_cast(width) * fmt.channels() * DataType(fmt.dtype()).size(); + return ImageDataStridedHost(fmt, MakeSinglePlaneBuffer(width, height, rowStride, basePtr)); +} + +// Single-plane GPU-resident Image wrapping a sentinel pointer via ImageWrapData. +// Use for batch tests where pushBack only reads the descriptor. +inline Image MakeFakeGpuImage(int32_t width, int32_t height, void* basePtr, ImageFormat fmt = FMT_RGB8) { + return ImageWrapData(MakeFakeHipData(width, height, basePtr, fmt)); +} + +// Host counterpart of MakeFakeGpuImage. +inline Image MakeFakeHostImage(int32_t width, int32_t height, void* basePtr, ImageFormat fmt = FMT_RGB8) { + return ImageWrapData(MakeFakeHostData(width, height, basePtr, fmt)); +} + +} // namespace tests +} // namespace roccv diff --git a/tests/roccv/cpp/include/test_helpers.hpp b/tests/roccv/cpp/include/test_helpers.hpp index 6c43053b..7ed56309 100644 --- a/tests/roccv/cpp/include/test_helpers.hpp +++ b/tests/roccv/cpp/include/test_helpers.hpp @@ -198,6 +198,12 @@ namespace tests { ". Expected no exceptions, but received the following exception: " + e.what()); \ } +// EXPECT_EQ pipes through std::to_string, so wrap enums/pointers/bools through +// these casts before comparing. +inline auto AsInt = [](auto v) { return static_cast(v); }; +inline auto AsAddr = [](const void* p) { return reinterpret_cast(p); }; +inline auto AsSize = [](auto v) { return static_cast(v); }; + /** * @brief Creates a NHWC tensor which contains data loaded from an image. * diff --git a/tests/roccv/cpp/src/tests/core/image/test_image.cpp b/tests/roccv/cpp/src/tests/core/image/test_image.cpp new file mode 100644 index 00000000..ce6ef69b --- /dev/null +++ b/tests/roccv/cpp/src/tests/core/image/test_image.cpp @@ -0,0 +1,375 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include + +#include +#include +#include +#include +#include + +#include "image_test_helpers.hpp" +#include "test_helpers.hpp" + +using namespace roccv; +using namespace roccv::tests; + +namespace { + +// ============================================================================= +// CalcRequirements +// ============================================================================= + +/** + * @brief Packed-row stride for a typical 3-channel uint8 image. Other fields + * propagate unchanged; remaining plane slots stay zeroed. + */ +void TestCalcRequirementsRgb8() { + auto reqs = Image::CalcRequirements({320, 240}, FMT_RGB8); + + EXPECT_EQ(reqs.size.w, 320); + EXPECT_EQ(reqs.size.h, 240); + EXPECT_EQ(reqs.format.channels(), 3); + EXPECT_EQ(reqs.planeRowStride[0], static_cast(320 * 3)); + EXPECT_EQ(reqs.planeRowStride[1], 0); + EXPECT_EQ(reqs.planeRowStride[5], 0); + EXPECT_EQ(reqs.alignBytes, 0); +} + +/** + * @brief Multi-byte dtype is reflected in the per-pixel byte count. + */ +void TestCalcRequirementsF32() { + auto reqs = Image::CalcRequirements({64, 64}, FMT_F32); + EXPECT_EQ(reqs.planeRowStride[0], static_cast(64 * 4)); +} + +/** + * @brief Single-channel U8 → row stride equals width. + */ +void TestCalcRequirementsU8() { + auto reqs = Image::CalcRequirements({100, 50}, FMT_U8); + EXPECT_EQ(reqs.planeRowStride[0], 100); +} + +/** + * @brief Width or height < 1 must throw INVALID_VALUE. + */ +void TestCalcRequirementsRejectsInvalidDims() { + EXPECT_EXCEPTION(Image::CalcRequirements({0, 100}, FMT_RGB8), eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(Image::CalcRequirements({100, 0}, FMT_RGB8), eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(Image::CalcRequirements({-5, 100}, FMT_RGB8), eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(Image::CalcRequirements({100, -5}, FMT_RGB8), eStatusType::INVALID_VALUE); +} + +/** + * @brief Large widths must not overflow during stride math; row stride must + * fit in int64. + */ +void TestCalcRequirementsLargeDims() { + // 8K image, RGBA8 (4 channels * 1 byte = 4 B/pixel) → 8192 * 4 = 32768 B/row. + auto reqs = Image::CalcRequirements({8192, 4320}, FMT_RGBA8); + EXPECT_EQ(reqs.planeRowStride[0], static_cast(8192 * 4)); +} + +// ============================================================================= +// Allocating constructors +// ============================================================================= + +/** + * @brief GPU-device ctor routes allocation through allocHipMem with the + * computed byte count. + */ +void TestImageHipAllocation() { + CountingAllocator alloc; + { + Image img({320, 240}, FMT_RGB8, alloc, eDeviceType::GPU); + + EXPECT_EQ(alloc.hipAllocs, 1); + EXPECT_EQ(alloc.hostAllocs, 0); + EXPECT_EQ(AsSize(alloc.lastAllocBytes), AsSize(320 * 3 * 240)); + + EXPECT_EQ(img.size().w, 320); + EXPECT_EQ(img.size().h, 240); + EXPECT_EQ(AsInt(img.device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(img.format().channels(), 3); + + // Image is still alive — buffer not yet freed. + EXPECT_EQ(alloc.hipFrees, 0); + } + // Image dropped — buffer freed exactly once via the matching allocator. + EXPECT_EQ(alloc.hipFrees, 1); +} + +/** + * @brief Same shape as the Hip test but for CPU residency. + */ +void TestImageHostAllocation() { + CountingAllocator alloc; + { + Image img({100, 50}, FMT_U8, alloc, eDeviceType::CPU); + + EXPECT_EQ(alloc.hostAllocs, 1); + EXPECT_EQ(alloc.hipAllocs, 0); + EXPECT_EQ(AsSize(alloc.lastAllocBytes), AsSize(100 * 50)); + EXPECT_EQ(AsInt(img.device()), AsInt(eDeviceType::CPU)); + } + EXPECT_EQ(alloc.hostFrees, 1); +} + +/** + * @brief Constructing from precomputed Requirements yields observably + * identical state to the (Size2D, ImageFormat) sugar form. + */ +void TestImageRequirementsCtor() { + CountingAllocator alloc; + auto reqs = Image::CalcRequirements({64, 32}, FMT_RGBA8); + + Image img(reqs, alloc, eDeviceType::GPU); + + EXPECT_EQ(img.size().w, 64); + EXPECT_EQ(img.size().h, 32); + EXPECT_EQ(img.format().channels(), 4); + EXPECT_EQ(AsSize(alloc.lastAllocBytes), AsSize(64 * 4 * 32)); +} + +// ============================================================================= +// Refcount / lifecycle +// ============================================================================= + +/** + * @brief Copying an Image bumps the refcount: both handles see the same + * underlying buffer, and free is deferred until the LAST handle drops. + */ +void TestImageCopySharesBuffer() { + CountingAllocator alloc; + void* buf = nullptr; + { + Image first({16, 16}, FMT_U8, alloc, eDeviceType::GPU); + buf = first.exportData().cast()->plane(0).basePtr; + + Image second = first; // refcount bump + EXPECT_EQ(alloc.hipAllocs, 1); // No new allocation. + EXPECT_EQ(AsAddr(second.exportData().cast()->plane(0).basePtr), AsAddr(buf)); + + // Drop `first`; buffer must NOT be freed yet — `second` still holds it. + { + Image sink = std::move(first); + } + EXPECT_EQ(alloc.hipFrees, 0); + } + // All handles dropped — exactly one free. + EXPECT_EQ(alloc.hipFrees, 1); +} + +/** + * @brief Move-construction transfers the buffer; the source is left empty. + * The buffer must still free exactly once (when the destination drops). + */ +void TestImageMoveSemantics() { + CountingAllocator alloc; + { + Image src({8, 8}, FMT_U8, alloc, eDeviceType::CPU); + void* srcBuf = src.exportData().cast()->plane(0).basePtr; + + Image dst = std::move(src); + EXPECT_EQ(AsAddr(dst.exportData().cast()->plane(0).basePtr), AsAddr(srcBuf)); + EXPECT_EQ(alloc.hostFrees, 0); + } + EXPECT_EQ(alloc.hostFrees, 1); +} + +// ============================================================================= +// exportData / exportData() +// ============================================================================= + +/** + * @brief exportData() returns an ImageData snapshot that mirrors the Image's + * size, format, device, and base pointer. + */ +void TestImageExportData() { + CountingAllocator alloc; + Image img({80, 60}, FMT_RGBA8, alloc, eDeviceType::GPU); + ImageData data = img.exportData(); + + EXPECT_EQ(AsInt(data.device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(data.format().channels(), 4); + + auto strided = data.cast(); + EXPECT_EQ(AsInt(strided.has_value()), 1); + EXPECT_EQ(strided->plane(0).width, 80); + EXPECT_EQ(strided->plane(0).height, 60); + EXPECT_EQ(strided->plane(0).rowStride, static_cast(80 * 4)); +} + +/** + * @brief Templated exportData() returns the matching subclass directly. + */ +void TestImageExportDataTypedSuccess() { + CountingAllocator alloc; + Image img({4, 4}, FMT_U8, alloc, eDeviceType::GPU); + + auto hip = img.exportData(); + EXPECT_EQ(AsInt(hip.device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(hip.plane(0).width, 4); +} + +/** + * @brief Templated exportData() throws std::bad_cast when the requested + * subclass does not match the underlying buffer kind. + */ +void TestImageExportDataTypedMismatch() { + CountingAllocator alloc; + Image img({4, 4}, FMT_U8, alloc, eDeviceType::GPU); + + bool threw = false; + try { + (void)img.exportData(); + } catch (const std::bad_cast&) { + threw = true; + } + EXPECT_EQ(AsInt(threw), 1); +} + +// ============================================================================= +// ImageWrapData +// ============================================================================= + +/** + * @brief View-only wrap (no cleanup callback) round-trips metadata and must + * not crash when the Image is destroyed (no free attempt on the sentinel ptr). + */ +void TestImageWrapDataViewOnly() { + Image wrapped = ImageWrapData(MakeFakeHipData(640, 480, FAKE_PTR_A)); + EXPECT_EQ(wrapped.size().w, 640); + EXPECT_EQ(wrapped.size().h, 480); + EXPECT_EQ(AsInt(wrapped.device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(AsAddr(wrapped.exportData().cast()->plane(0).basePtr), AsAddr(FAKE_PTR_A)); +} + +/** + * @brief Wrap with a cleanup callback: the callback fires exactly once when + * the last Image handle goes out of scope. + */ +void TestImageWrapDataCleanupFires() { + int callbackInvocations = 0; + { + Image wrapped = + ImageWrapData(MakeFakeHipData(100, 100, FAKE_PTR_A), [&](const ImageData&) { ++callbackInvocations; }); + EXPECT_EQ(callbackInvocations, 0); // Not fired during normal use. + } + EXPECT_EQ(callbackInvocations, 1); +} + +/** + * @brief Cleanup callback receives the original wrapped ImageData snapshot — + * the captured base pointer must match what was passed to ImageWrapData. + */ +void TestImageWrapDataCleanupReceivesData() { + void* receivedBasePtr = nullptr; + { + Image wrapped = ImageWrapData(MakeFakeHipData(50, 50, FAKE_PTR_A), [&](const ImageData& d) { + receivedBasePtr = d.cast()->plane(0).basePtr; + }); + } + EXPECT_EQ(AsAddr(receivedBasePtr), AsAddr(FAKE_PTR_A)); +} + +/** + * @brief Cleanup must fire only on LAST handle drop — copies bump the + * refcount, intermediate drops do nothing. + */ +void TestImageWrapDataCleanupFiresOnce() { + int callbackInvocations = 0; + { + Image first = + ImageWrapData(MakeFakeHipData(10, 10, FAKE_PTR_A), [&](const ImageData&) { ++callbackInvocations; }); + Image second = first; // refcount = 2 + Image third = first; // refcount = 3 + { + Image fourth = third; + (void)fourth; + } // dropped → refcount = 3 + EXPECT_EQ(callbackInvocations, 0); + // first, second, third still alive at scope exit + } + EXPECT_EQ(callbackInvocations, 1); +} + +/** + * @brief Wrapped Image's accessors mirror the wrapped ImageData verbatim — + * size, format, device, and base pointer all round-trip unchanged. + */ +void TestImageWrapDataAccessors() { + auto fake = MakeFakeHipData(123, 45, FAKE_PTR_A, FMT_RGBA8); + Image wrapped = ImageWrapData(fake); + + EXPECT_EQ(wrapped.size().w, 123); + EXPECT_EQ(wrapped.size().h, 45); + EXPECT_EQ(wrapped.format().channels(), 4); + EXPECT_EQ(AsInt(wrapped.device()), AsInt(eDeviceType::GPU)); + + auto strided = wrapped.exportData().cast(); + EXPECT_EQ(AsInt(strided.has_value()), 1); + EXPECT_EQ(strided->plane(0).width, 123); + EXPECT_EQ(strided->plane(0).height, 45); + EXPECT_EQ(AsAddr(strided->plane(0).basePtr), AsAddr(FAKE_PTR_A)); +} + +} // namespace + +int main(int argc, char** argv) { + (void)argc; + (void)argv; + TEST_CASES_BEGIN(); + + // CalcRequirements + TEST_CASE(TestCalcRequirementsRgb8()); + TEST_CASE(TestCalcRequirementsF32()); + TEST_CASE(TestCalcRequirementsU8()); + TEST_CASE(TestCalcRequirementsRejectsInvalidDims()); + TEST_CASE(TestCalcRequirementsLargeDims()); + + // Allocating constructors + TEST_CASE(TestImageHipAllocation()); + TEST_CASE(TestImageHostAllocation()); + TEST_CASE(TestImageRequirementsCtor()); + + // Refcount / lifecycle + TEST_CASE(TestImageCopySharesBuffer()); + TEST_CASE(TestImageMoveSemantics()); + + // exportData + TEST_CASE(TestImageExportData()); + TEST_CASE(TestImageExportDataTypedSuccess()); + TEST_CASE(TestImageExportDataTypedMismatch()); + + // ImageWrapData + TEST_CASE(TestImageWrapDataViewOnly()); + TEST_CASE(TestImageWrapDataCleanupFires()); + TEST_CASE(TestImageWrapDataCleanupReceivesData()); + TEST_CASE(TestImageWrapDataCleanupFiresOnce()); + TEST_CASE(TestImageWrapDataAccessors()); + + TEST_CASES_END(); +} diff --git a/tests/roccv/cpp/src/tests/core/image/test_image_batch_data.cpp b/tests/roccv/cpp/src/tests/core/image/test_image_batch_data.cpp new file mode 100644 index 00000000..31e449e6 --- /dev/null +++ b/tests/roccv/cpp/src/tests/core/image/test_image_batch_data.cpp @@ -0,0 +1,283 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include + +#include +#include +#include +#include + +#include "image_test_helpers.hpp" +#include "test_helpers.hpp" + +using namespace roccv; +using namespace roccv::tests; + +namespace { + +// Static descriptor/format storage for the batch buffer. These are real host +// allocations (so the pointers are valid) but the batch tests only read +// metadata back out of them; nothing dereferences the per-image basePtr fields. +ImageBufferStrided g_imageList[2]; +ImageFormat g_formatList[2] = {FMT_RGB8, FMT_RGB8}; +ImageFormat g_hostFormatList[2] = {FMT_RGB8, FMT_RGB8}; + +// Builds a homogeneous two-image varshape descriptor with a known bounding box +// and uniqueFormat. The returned struct's pointers reference module-static +// arrays so addresses remain stable across calls within a test. +ImageBatchVarShapeBufferStrided MakeHomogeneousBuffer() { + g_imageList[0] = MakeSinglePlaneBuffer(640, 480, 640 * 3, FAKE_PTR_A); + g_imageList[1] = MakeSinglePlaneBuffer(320, 240, 320 * 3, FAKE_PTR_B); + g_formatList[0] = FMT_RGB8; + g_formatList[1] = FMT_RGB8; + g_hostFormatList[0] = FMT_RGB8; + g_hostFormatList[1] = FMT_RGB8; + + ImageBatchVarShapeBufferStrided buf{}; + buf.uniqueFormat = FMT_RGB8; + buf.maxWidth = 640; + buf.maxHeight = 480; + buf.formatList = g_formatList; + buf.hostFormatList = g_hostFormatList; + buf.imageList = g_imageList; + return buf; +} + +/** + * @brief Verifies HIP-strided varshape construction populates all observable + * state and tags itself as GPU-resident. + */ +void TestImageBatchVarShapeDataStridedHipConstruction() { + auto buf = MakeHomogeneousBuffer(); + ImageBatchVarShapeDataStridedHip data(2, buf); + + EXPECT_EQ(AsInt(data.device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(data.numImages(), 2); + EXPECT_EQ(data.maxSize().w, 640); + EXPECT_EQ(data.maxSize().h, 480); + EXPECT_EQ(data.uniqueFormat().channels(), 3); + EXPECT_EQ(AsAddr(data.formatList()), AsAddr(g_formatList)); + EXPECT_EQ(AsAddr(data.hostFormatList()), AsAddr(g_hostFormatList)); + EXPECT_EQ(AsAddr(data.imageList()), AsAddr(g_imageList)); + EXPECT_EQ(data.imageList()[0].planes[0].width, 640); + EXPECT_EQ(data.imageList()[1].planes[0].width, 320); +} + +/** + * @brief Same shape as the Hip test but for Host-resident varshape data. + */ +void TestImageBatchVarShapeDataStridedHostConstruction() { + auto buf = MakeHomogeneousBuffer(); + ImageBatchVarShapeDataStridedHost data(2, buf); + + EXPECT_EQ(AsInt(data.device()), AsInt(eDeviceType::CPU)); + EXPECT_EQ(data.numImages(), 2); + EXPECT_EQ(data.maxSize().w, 640); + EXPECT_EQ(data.maxSize().h, 480); + EXPECT_EQ(data.uniqueFormat().channels(), 3); + EXPECT_EQ(AsAddr(data.imageList()), AsAddr(g_imageList)); +} + +/** + * @brief Empty batch: maxSize collapses to 0x0 and uniqueFormat is FMT_NONE. + * Producers signal "no images" via numImages == 0; the buffer fields stay + * valid pointers but get ignored. + */ +void TestImageBatchVarShapeDataEmpty() { + ImageBatchVarShapeBufferStrided buf{}; + buf.uniqueFormat = FMT_NONE; + buf.maxWidth = 0; + buf.maxHeight = 0; + buf.formatList = g_formatList; + buf.hostFormatList = g_hostFormatList; + buf.imageList = g_imageList; + + ImageBatchVarShapeDataStridedHip data(0, buf); + + EXPECT_EQ(data.numImages(), 0); + EXPECT_EQ(data.maxSize().w, 0); + EXPECT_EQ(data.maxSize().h, 0); + EXPECT_EQ(AsInt(data.uniqueFormat() == FMT_NONE), 1); +} + +/** + * @brief Heterogeneous formats: per-image formatList carries each entry + * verbatim; uniqueFormat is FMT_NONE since no single format spans the batch. + */ +void TestImageBatchVarShapeDataHeterogeneousFormats() { + g_imageList[0] = MakeSinglePlaneBuffer(640, 480, 640 * 3, FAKE_PTR_A); + g_imageList[1] = MakeSinglePlaneBuffer(320, 240, 320 * 4, FAKE_PTR_B); + g_formatList[0] = FMT_RGB8; + g_formatList[1] = FMT_RGBA8; + g_hostFormatList[0] = FMT_RGB8; + g_hostFormatList[1] = FMT_RGBA8; + + ImageBatchVarShapeBufferStrided buf{}; + buf.uniqueFormat = FMT_NONE; + buf.maxWidth = 640; + buf.maxHeight = 480; + buf.formatList = g_formatList; + buf.hostFormatList = g_hostFormatList; + buf.imageList = g_imageList; + + ImageBatchVarShapeDataStridedHip data(2, buf); + + EXPECT_EQ(AsInt(data.uniqueFormat() == FMT_NONE), 1); + EXPECT_EQ(AsInt(data.hostFormatList()[0] == FMT_RGB8), 1); + EXPECT_EQ(AsInt(data.hostFormatList()[1] == FMT_RGBA8), 1); +} + +/** + * @brief The two leaf ctors (taking ImageBatchBuffer vs the concrete strided + * buffer directly) must produce observably identical state. + */ +void TestImageBatchVarShapeDataSugarCtor() { + auto buf = MakeHomogeneousBuffer(); + + ImageBatchVarShapeDataStridedHip wide(2, ImageBatchBuffer{.varShapeStrided = buf}); + ImageBatchVarShapeDataStridedHip sugar(2, buf); + + EXPECT_EQ(AsInt(wide.device()), AsInt(sugar.device())); + EXPECT_EQ(wide.numImages(), sugar.numImages()); + EXPECT_EQ(wide.maxSize().w, sugar.maxSize().w); + EXPECT_EQ(wide.maxSize().h, sugar.maxSize().h); + EXPECT_EQ(AsAddr(wide.imageList()), AsAddr(sugar.imageList())); + + ImageBatchVarShapeDataStridedHost wideHost(2, ImageBatchBuffer{.varShapeStrided = buf}); + ImageBatchVarShapeDataStridedHost sugarHost(2, buf); + EXPECT_EQ(AsInt(wideHost.device()), AsInt(sugarHost.device())); + EXPECT_EQ(AsAddr(wideHost.imageList()), AsAddr(sugarHost.imageList())); +} + +/** + * @brief IsCompatibleKind on each level discriminates the buffer kinds it + * accepts. Base accepts anything-but-NONE; VarShape and VarShapeStrided accept + * both Hip and Host varshape; leaves accept only their own. + */ +void TestImageBatchDataIsCompatibleKind() { + EXPECT_EQ(AsInt(ImageBatchData::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_BUFFER_NONE)), 0); + EXPECT_EQ(AsInt(ImageBatchData::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP)), + 1); + EXPECT_EQ(AsInt(ImageBatchData::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST)), + 1); + + EXPECT_EQ(AsInt(ImageBatchVarShapeData::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_BUFFER_NONE)), 0); + EXPECT_EQ( + AsInt(ImageBatchVarShapeData::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP)), + 1); + EXPECT_EQ( + AsInt(ImageBatchVarShapeData::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST)), + 1); + + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStrided::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_BUFFER_NONE)), 0); + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStrided::IsCompatibleKind( + ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP)), + 1); + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStrided::IsCompatibleKind( + ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST)), + 1); + + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStridedHip::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_BUFFER_NONE)), + 0); + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStridedHip::IsCompatibleKind( + ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP)), + 1); + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStridedHip::IsCompatibleKind( + ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST)), + 0); + + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStridedHost::IsCompatibleKind(ImageBatchBufferType::IMAGE_BATCH_BUFFER_NONE)), + 0); + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStridedHost::IsCompatibleKind( + ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HIP)), + 0); + EXPECT_EQ(AsInt(ImageBatchVarShapeDataStridedHost::IsCompatibleKind( + ImageBatchBufferType::IMAGE_BATCH_VARSHAPE_BUFFER_STRIDED_HOST)), + 1); +} + +/** + * @brief Round-trip a derived ImageBatchData through the base reference and + * back via cast<>(). Successful casts must preserve every observable field; + * casts to incompatible kinds must return std::nullopt. + */ +void TestImageBatchDataCast() { + auto buf = MakeHomogeneousBuffer(); + + // Hip → base → Hip should round-trip; intermediate VarShape/Strided also + // succeed; Hip → Host fails. + { + ImageBatchVarShapeDataStridedHip hip(2, buf); + const ImageBatchData& base = hip; + + auto asHip = base.cast(); + EXPECT_EQ(AsInt(asHip.has_value()), 1); + EXPECT_EQ(AsInt(asHip->device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(asHip->numImages(), 2); + EXPECT_EQ(asHip->maxSize().w, 640); + EXPECT_EQ(AsAddr(asHip->imageList()), AsAddr(g_imageList)); + + auto asStrided = base.cast(); + EXPECT_EQ(AsInt(asStrided.has_value()), 1); + EXPECT_EQ(AsInt(asStrided->device()), AsInt(eDeviceType::GPU)); + + auto asVar = base.cast(); + EXPECT_EQ(AsInt(asVar.has_value()), 1); + EXPECT_EQ(asVar->maxSize().h, 480); + + auto asHost = base.cast(); + EXPECT_EQ(AsInt(asHost.has_value()), 0); + } + + // Symmetrically: Host → base → Host succeeds, Host → Hip fails. + { + ImageBatchVarShapeDataStridedHost host(2, buf); + const ImageBatchData& base = host; + + auto asHost = base.cast(); + EXPECT_EQ(AsInt(asHost.has_value()), 1); + EXPECT_EQ(AsInt(asHost->device()), AsInt(eDeviceType::CPU)); + EXPECT_EQ(asHost->numImages(), 2); + + auto asHip = base.cast(); + EXPECT_EQ(AsInt(asHip.has_value()), 0); + } +} + +} // namespace + +int main(int argc, char** argv) { + (void)argc; + (void)argv; + TEST_CASES_BEGIN(); + + TEST_CASE(TestImageBatchVarShapeDataStridedHipConstruction()); + TEST_CASE(TestImageBatchVarShapeDataStridedHostConstruction()); + TEST_CASE(TestImageBatchVarShapeDataEmpty()); + TEST_CASE(TestImageBatchVarShapeDataHeterogeneousFormats()); + TEST_CASE(TestImageBatchVarShapeDataSugarCtor()); + TEST_CASE(TestImageBatchDataIsCompatibleKind()); + TEST_CASE(TestImageBatchDataCast()); + + TEST_CASES_END(); +} diff --git a/tests/roccv/cpp/src/tests/core/image/test_image_batch_var_shape.cpp b/tests/roccv/cpp/src/tests/core/image/test_image_batch_var_shape.cpp new file mode 100644 index 00000000..ec148240 --- /dev/null +++ b/tests/roccv/cpp/src/tests/core/image/test_image_batch_var_shape.cpp @@ -0,0 +1,394 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "image_test_helpers.hpp" +#include "test_helpers.hpp" + +using namespace roccv; +using namespace roccv::tests; + +namespace { + +// ============================================================================= +// Construction +// ============================================================================= + +void TestConstruction() { + CountingAllocator alloc; + { + ImageBatchVarShape batch(8, alloc); + EXPECT_EQ(batch.capacity(), 8); + EXPECT_EQ(batch.numImages(), 0); + EXPECT_EQ(AsInt(batch.begin() == batch.end()), 1); + } + EXPECT_EQ(alloc.hipAllocs, 2); + EXPECT_EQ(alloc.pinnedAllocs, 2); + EXPECT_EQ(alloc.hipFrees, 2); + EXPECT_EQ(alloc.pinnedFrees, 2); +} + +void TestConstructionRejectsBadCapacity() { + CountingAllocator alloc; + EXPECT_EXCEPTION(ImageBatchVarShape(0, alloc), eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(ImageBatchVarShape(-3, alloc), eStatusType::INVALID_VALUE); +} + +// ============================================================================= +// pushBack — basic +// ============================================================================= + +void TestPushBackSingle() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + + Image img = MakeFakeGpuImage(640, 480, FAKE_PTR_A); + batch.pushBack(img); + + EXPECT_EQ(batch.numImages(), 1); + EXPECT_EQ(batch[0].size().w, 640); + EXPECT_EQ(batch[0].size().h, 480); + EXPECT_EQ(AsInt(batch[0].format() == FMT_RGB8), 1); +} + +void TestPushBackMultipleHeterogeneousSizes() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + + batch.pushBack(MakeFakeGpuImage(640, 480, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(320, 240, FAKE_PTR_B)); + batch.pushBack(MakeFakeGpuImage(800, 200, FAKE_PTR_C)); + + EXPECT_EQ(batch.numImages(), 3); + EXPECT_EQ(batch.maxSize().w, 800); + EXPECT_EQ(batch.maxSize().h, 480); + EXPECT_EQ(AsInt(batch.uniqueFormat() == FMT_RGB8), 1); +} + +void TestPushBackIteratorRange() { + CountingAllocator alloc; + ImageBatchVarShape batch(8, alloc); + + std::vector imgs; + imgs.push_back(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + imgs.push_back(MakeFakeGpuImage(200, 200, FAKE_PTR_B)); + imgs.push_back(MakeFakeGpuImage(300, 300, FAKE_PTR_C)); + + batch.pushBack(imgs.begin(), imgs.end()); + + EXPECT_EQ(batch.numImages(), 3); + EXPECT_EQ(batch.maxSize().w, 300); +} + +// ============================================================================= +// pushBack — validation +// ============================================================================= + +void TestPushBackCapacityOverflow() { + CountingAllocator alloc; + ImageBatchVarShape batch(2, alloc); + + batch.pushBack(MakeFakeGpuImage(64, 64, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(64, 64, FAKE_PTR_B)); + + EXPECT_EXCEPTION(batch.pushBack(MakeFakeGpuImage(64, 64, FAKE_PTR_C)), eStatusType::OUT_OF_BOUNDS); +} + +void TestPushBackHostImageRejected() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + + Image cpuImg = MakeFakeHostImage(64, 64, FAKE_PTR_A, FMT_U8); + EXPECT_EXCEPTION(batch.pushBack(cpuImg), eStatusType::INVALID_VALUE); +} + +// Note: pushBack's single-plane validation is defense-in-depth — Image's own +// exportData() (image.cpp:118) currently hardcodes numPlanes=1 regardless of +// the underlying buffer, so the public API can't construct a multi-plane Image +// for this guard to fire on. The test would need to be revisited when planar +// formats land in Image itself. + +void TestPushBackRangeRollbackOnFailure() { + CountingAllocator alloc; + ImageBatchVarShape batch(8, alloc); + + // Pre-populate so we can confirm the rollback restores exactly the + // pre-call state, not just back to zero. + batch.pushBack(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + EXPECT_EQ(batch.numImages(), 1); + + // Mid-range CPU image — should rollback the partially-pushed entries. + std::vector imgs; + imgs.push_back(MakeFakeGpuImage(200, 200, FAKE_PTR_B)); + imgs.push_back(MakeFakeHostImage(300, 300, FAKE_PTR_C)); // Will throw. + + EXPECT_EXCEPTION(batch.pushBack(imgs.begin(), imgs.end()), eStatusType::INVALID_VALUE); + + // Pre-call state is intact: 1 image, original maxSize. + EXPECT_EQ(batch.numImages(), 1); + EXPECT_EQ(batch.maxSize().w, 100); +} + +void TestPushBackRangeOverflowPrechecked() { + CountingAllocator alloc; + ImageBatchVarShape batch(2, alloc); + + std::vector imgs; + imgs.push_back(MakeFakeGpuImage(10, 10, FAKE_PTR_A)); + imgs.push_back(MakeFakeGpuImage(20, 20, FAKE_PTR_B)); + imgs.push_back(MakeFakeGpuImage(30, 30, FAKE_PTR_C)); // 3rd overflows capacity 2. + + EXPECT_EXCEPTION(batch.pushBack(imgs.begin(), imgs.end()), eStatusType::OUT_OF_BOUNDS); + // Pre-checked: nothing was pushed. + EXPECT_EQ(batch.numImages(), 0); +} + +// ============================================================================= +// popBack / clear +// ============================================================================= + +void TestPopBack() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + + batch.pushBack(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(200, 200, FAKE_PTR_B)); + batch.popBack(); + + EXPECT_EQ(batch.numImages(), 1); + // maxSize was reset on pop; the rescan should drop back to 100. + EXPECT_EQ(batch.maxSize().w, 100); +} + +void TestPopBackMultiple() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + + batch.pushBack(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(200, 200, FAKE_PTR_B)); + batch.pushBack(MakeFakeGpuImage(300, 300, FAKE_PTR_C)); + batch.popBack(2); + + EXPECT_EQ(batch.numImages(), 1); + EXPECT_EQ(batch.maxSize().w, 100); +} + +void TestPopBackUnderflow() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + batch.pushBack(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + + EXPECT_EXCEPTION(batch.popBack(2), eStatusType::OUT_OF_BOUNDS); + // State preserved. + EXPECT_EQ(batch.numImages(), 1); +} + +void TestClearAndReuse() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + + batch.pushBack(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(200, 200, FAKE_PTR_B)); + batch.clear(); + + EXPECT_EQ(batch.numImages(), 0); + EXPECT_EQ(batch.maxSize().w, 0); + EXPECT_EQ(AsInt(batch.uniqueFormat() == FMT_NONE), 1); + + // Reuse after clear. + batch.pushBack(MakeFakeGpuImage(50, 50, FAKE_PTR_C, FMT_U8)); + EXPECT_EQ(batch.numImages(), 1); + EXPECT_EQ(AsInt(batch.uniqueFormat() == FMT_U8), 1); +} + +// ============================================================================= +// uniqueFormat / maxSize cache +// ============================================================================= + +void TestUniqueFormatHomogeneous() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + batch.pushBack(MakeFakeGpuImage(64, 64, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(128, 128, FAKE_PTR_B)); + EXPECT_EQ(AsInt(batch.uniqueFormat() == FMT_RGB8), 1); +} + +void TestUniqueFormatHeterogeneous() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + batch.pushBack(MakeFakeGpuImage(64, 64, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(64, 64, FAKE_PTR_B, FMT_RGBA8)); + EXPECT_EQ(AsInt(batch.uniqueFormat() == FMT_NONE), 1); +} + +void TestUniqueFormatEmptyBatch() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + EXPECT_EQ(AsInt(batch.uniqueFormat() == FMT_NONE), 1); + EXPECT_EQ(batch.maxSize().w, 0); + EXPECT_EQ(batch.maxSize().h, 0); +} + +// ============================================================================= +// exportData +// ============================================================================= + +// exportData tests use the default allocator instead of CountingAllocator +// because they exercise the real H2D hipMemcpyAsync, which requires the +// device-side buffer to be a real hipMalloc'd pointer. + +void TestExportDataEmpty() { + ImageBatchVarShape batch(4); + + auto data = batch.exportData(0); + EXPECT_EQ(data.numImages(), 0); + EXPECT_EQ(data.maxSize().w, 0); + EXPECT_EQ(data.maxSize().h, 0); + EXPECT_EQ(AsInt(data.uniqueFormat() == FMT_NONE), 1); + EXPECT_EQ(AsInt(data.device()), AsInt(eDeviceType::GPU)); +} + +void TestExportDataMetadata() { + ImageBatchVarShape batch(4); + batch.pushBack(MakeFakeGpuImage(640, 480, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(320, 240, FAKE_PTR_B)); + + auto data = batch.exportData(0); + EXPECT_EQ(data.numImages(), 2); + EXPECT_EQ(data.maxSize().w, 640); + EXPECT_EQ(data.maxSize().h, 480); + EXPECT_EQ(AsInt(data.uniqueFormat() == FMT_RGB8), 1); + EXPECT_EQ(AsInt(data.imageList() != nullptr), 1); + EXPECT_EQ(AsInt(data.formatList() != nullptr), 1); + EXPECT_EQ(AsInt(data.hostFormatList() != nullptr), 1); + // Pinned host mirror format entries are immediately host-readable. + EXPECT_EQ(AsInt(data.hostFormatList()[0] == FMT_RGB8), 1); + EXPECT_EQ(AsInt(data.hostFormatList()[1] == FMT_RGB8), 1); +} + +void TestExportDataCastRoundTrip() { + ImageBatchVarShape batch(4); + batch.pushBack(MakeFakeGpuImage(64, 64, FAKE_PTR_A)); + + auto hipData = batch.exportData(0); + EXPECT_EQ(hipData.numImages(), 1); + EXPECT_EQ(AsInt(hipData.device()), AsInt(eDeviceType::GPU)); + + // Cast through the base reference: succeeds for compatible kinds, nullopt + // for the host-resident leaf. + const ImageBatchData& base = hipData; + EXPECT_EQ(AsInt(base.cast().has_value()), 1); + EXPECT_EQ(AsInt(base.cast().has_value()), 0); +} + +// ============================================================================= +// Move semantics +// ============================================================================= + +void TestMoveConstruction() { + CountingAllocator alloc; + { + ImageBatchVarShape src(4, alloc); + src.pushBack(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + src.pushBack(MakeFakeGpuImage(200, 200, FAKE_PTR_B)); + + ImageBatchVarShape dst(std::move(src)); + EXPECT_EQ(dst.numImages(), 2); + EXPECT_EQ(dst.maxSize().w, 200); + + // Source is valid-but-empty; destructor must not double-free. + EXPECT_EQ(src.numImages(), 0); + EXPECT_EQ(src.capacity(), 0); + } + // Exactly one set of allocations should have been freed. + EXPECT_EQ(alloc.hipAllocs, alloc.hipFrees); + EXPECT_EQ(alloc.pinnedAllocs, alloc.pinnedFrees); +} + +// ============================================================================= +// Iterator +// ============================================================================= + +void TestIteratorRangeFor() { + CountingAllocator alloc; + ImageBatchVarShape batch(4, alloc); + batch.pushBack(MakeFakeGpuImage(100, 100, FAKE_PTR_A)); + batch.pushBack(MakeFakeGpuImage(200, 200, FAKE_PTR_B)); + batch.pushBack(MakeFakeGpuImage(300, 300, FAKE_PTR_C)); + + int32_t expectedW = 100; + int32_t count = 0; + for (const Image& img : batch) { + EXPECT_EQ(img.size().w, expectedW); + expectedW += 100; + ++count; + } + EXPECT_EQ(count, 3); +} + +} // namespace + +int main(int argc, char** argv) { + (void)argc; + (void)argv; + TEST_CASES_BEGIN(); + + TEST_CASE(TestConstruction()); + TEST_CASE(TestConstructionRejectsBadCapacity()); + + TEST_CASE(TestPushBackSingle()); + TEST_CASE(TestPushBackMultipleHeterogeneousSizes()); + TEST_CASE(TestPushBackIteratorRange()); + + TEST_CASE(TestPushBackCapacityOverflow()); + TEST_CASE(TestPushBackHostImageRejected()); + TEST_CASE(TestPushBackRangeRollbackOnFailure()); + TEST_CASE(TestPushBackRangeOverflowPrechecked()); + + TEST_CASE(TestPopBack()); + TEST_CASE(TestPopBackMultiple()); + TEST_CASE(TestPopBackUnderflow()); + TEST_CASE(TestClearAndReuse()); + + TEST_CASE(TestUniqueFormatHomogeneous()); + TEST_CASE(TestUniqueFormatHeterogeneous()); + TEST_CASE(TestUniqueFormatEmptyBatch()); + + TEST_CASE(TestExportDataEmpty()); + TEST_CASE(TestExportDataMetadata()); + TEST_CASE(TestExportDataCastRoundTrip()); + + TEST_CASE(TestMoveConstruction()); + + TEST_CASE(TestIteratorRangeFor()); + + TEST_CASES_END(); +} diff --git a/tests/roccv/cpp/src/tests/core/image/test_image_data.cpp b/tests/roccv/cpp/src/tests/core/image/test_image_data.cpp new file mode 100644 index 00000000..8a7945fd --- /dev/null +++ b/tests/roccv/cpp/src/tests/core/image/test_image_data.cpp @@ -0,0 +1,207 @@ +/* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include + +#include +#include +#include + +#include "image_test_helpers.hpp" +#include "test_helpers.hpp" + +using namespace roccv; +using namespace roccv::tests; + +namespace { + +ImageBufferStrided MakeThreePlaneBuffer() { + // Mimics a planar layout (e.g. YUV420-style) with sub-sampled chroma — three + // planes of differing dimensions and strides backed by distinct buffers. + ImageBufferStrided buf{}; + buf.numPlanes = 3; + buf.planes[0] = {1920, 1080, 1920, FAKE_PTR_A}; // Y full-resolution + buf.planes[1] = {960, 540, 960, FAKE_PTR_B}; // U sub-sampled + buf.planes[2] = {960, 540, 960, FAKE_PTR_C}; // V sub-sampled + return buf; +} + +/** + * @brief Verifies HIP-strided construction populates all observable state and + * tags itself as GPU-resident. + */ +void TestImageDataStridedHipConstruction() { + auto buf = MakeSinglePlaneBuffer(640, 480, 640 * 3, FAKE_PTR_A); + ImageDataStridedHip data(FMT_RGB8, buf); + + EXPECT_EQ(AsInt(data.device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(data.numPlanes(), 1); + EXPECT_EQ(data.size().w, 640); + EXPECT_EQ(data.size().h, 480); + EXPECT_EQ(data.plane(0).width, 640); + EXPECT_EQ(data.plane(0).height, 480); + EXPECT_EQ(data.plane(0).rowStride, static_cast(640 * 3)); + EXPECT_EQ(AsAddr(data.plane(0).basePtr), AsAddr(FAKE_PTR_A)); + EXPECT_EQ(data.format().channels(), 3); +} + +/** + * @brief Same shape as the Hip test but for Host-resident strided data. + */ +void TestImageDataStridedHostConstruction() { + auto buf = MakeSinglePlaneBuffer(320, 240, 320, FAKE_PTR_B); + ImageDataStridedHost data(FMT_U8, buf); + + EXPECT_EQ(AsInt(data.device()), AsInt(eDeviceType::CPU)); + EXPECT_EQ(data.numPlanes(), 1); + EXPECT_EQ(data.size().w, 320); + EXPECT_EQ(data.size().h, 240); + EXPECT_EQ(AsAddr(data.plane(0).basePtr), AsAddr(FAKE_PTR_B)); + EXPECT_EQ(data.format().channels(), 1); +} + +/** + * @brief Multi-plane buffers must round-trip per-plane dimensions and pointers + * unchanged. size() reports plane 0 by convention; planes 1..N may be smaller. + */ +void TestImageDataStridedMultiPlane() { + auto buf = MakeThreePlaneBuffer(); + ImageDataStridedHip data(FMT_U8, buf); + + EXPECT_EQ(data.numPlanes(), 3); + EXPECT_EQ(data.size().w, 1920); + EXPECT_EQ(data.size().h, 1080); + + EXPECT_EQ(data.plane(0).width, 1920); + EXPECT_EQ(data.plane(0).height, 1080); + EXPECT_EQ(AsAddr(data.plane(0).basePtr), AsAddr(FAKE_PTR_A)); + + EXPECT_EQ(data.plane(1).width, 960); + EXPECT_EQ(data.plane(1).height, 540); + EXPECT_EQ(AsAddr(data.plane(1).basePtr), AsAddr(FAKE_PTR_B)); + + EXPECT_EQ(data.plane(2).width, 960); + EXPECT_EQ(data.plane(2).height, 540); + EXPECT_EQ(AsAddr(data.plane(2).basePtr), AsAddr(FAKE_PTR_C)); +} + +/** + * @brief The two leaf ctors (taking ImageBuffer vs ImageBufferStrided directly) + * must produce observably identical state. + */ +void TestImageDataStridedSugarCtor() { + auto buf = MakeSinglePlaneBuffer(100, 200, 400, FAKE_PTR_A); + + ImageDataStridedHip wide(FMT_RGBA8, ImageBuffer{.strided = buf}); + ImageDataStridedHip sugar(FMT_RGBA8, buf); + + EXPECT_EQ(AsInt(wide.device()), AsInt(sugar.device())); + EXPECT_EQ(wide.numPlanes(), sugar.numPlanes()); + EXPECT_EQ(AsAddr(wide.plane(0).basePtr), AsAddr(sugar.plane(0).basePtr)); + EXPECT_EQ(wide.plane(0).rowStride, sugar.plane(0).rowStride); + + ImageDataStridedHost wideHost(FMT_U8, ImageBuffer{.strided = buf}); + ImageDataStridedHost sugarHost(FMT_U8, buf); + EXPECT_EQ(AsInt(wideHost.device()), AsInt(sugarHost.device())); + EXPECT_EQ(AsAddr(wideHost.plane(0).basePtr), AsAddr(sugarHost.plane(0).basePtr)); +} + +/** + * @brief IsCompatibleKind on each level discriminates the buffer kinds it + * accepts. Base accepts anything-but-NONE; Strided accepts both Hip and Host; + * leaves accept only their own. + */ +void TestImageDataIsCompatibleKind() { + EXPECT_EQ(AsInt(ImageData::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_NONE)), 0); + EXPECT_EQ(AsInt(ImageData::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HIP)), 1); + EXPECT_EQ(AsInt(ImageData::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HOST)), 1); + + EXPECT_EQ(AsInt(ImageDataStrided::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_NONE)), 0); + EXPECT_EQ(AsInt(ImageDataStrided::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HIP)), 1); + EXPECT_EQ(AsInt(ImageDataStrided::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HOST)), 1); + + EXPECT_EQ(AsInt(ImageDataStridedHip::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_NONE)), 0); + EXPECT_EQ(AsInt(ImageDataStridedHip::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HIP)), 1); + EXPECT_EQ(AsInt(ImageDataStridedHip::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HOST)), 0); + + EXPECT_EQ(AsInt(ImageDataStridedHost::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_NONE)), 0); + EXPECT_EQ(AsInt(ImageDataStridedHost::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HIP)), 0); + EXPECT_EQ(AsInt(ImageDataStridedHost::IsCompatibleKind(ImageBufferType::IMAGE_BUFFER_STRIDED_HOST)), 1); +} + +/** + * @brief Round-trip a derived ImageData through the base reference and back + * via cast<>(). Successful casts must preserve every observable field; casts + * to incompatible kinds must return std::nullopt. + */ +void TestImageDataCast() { + auto buf = MakeSinglePlaneBuffer(800, 600, 800 * 4, FAKE_PTR_A); + + // Hip → base → Hip should round-trip, Hip → Host should fail. + { + ImageDataStridedHip hip(FMT_RGBA8, buf); + const ImageData& base = hip; + + auto asHip = base.cast(); + EXPECT_EQ(AsInt(asHip.has_value()), 1); + EXPECT_EQ(AsInt(asHip->device()), AsInt(eDeviceType::GPU)); + EXPECT_EQ(AsAddr(asHip->plane(0).basePtr), AsAddr(FAKE_PTR_A)); + EXPECT_EQ(asHip->plane(0).width, 800); + + auto asStrided = base.cast(); + EXPECT_EQ(AsInt(asStrided.has_value()), 1); + EXPECT_EQ(AsInt(asStrided->device()), AsInt(eDeviceType::GPU)); + + auto asHost = base.cast(); + EXPECT_EQ(AsInt(asHost.has_value()), 0); + } + + // Symmetrically: Host → base → Host succeeds, Host → Hip fails. + { + ImageDataStridedHost host(FMT_RGBA8, buf); + const ImageData& base = host; + + auto asHost = base.cast(); + EXPECT_EQ(AsInt(asHost.has_value()), 1); + EXPECT_EQ(AsInt(asHost->device()), AsInt(eDeviceType::CPU)); + + auto asHip = base.cast(); + EXPECT_EQ(AsInt(asHip.has_value()), 0); + } +} + +} // namespace + +int main(int argc, char** argv) { + (void)argc; + (void)argv; + TEST_CASES_BEGIN(); + + TEST_CASE(TestImageDataStridedHipConstruction()); + TEST_CASE(TestImageDataStridedHostConstruction()); + TEST_CASE(TestImageDataStridedMultiPlane()); + TEST_CASE(TestImageDataStridedSugarCtor()); + TEST_CASE(TestImageDataIsCompatibleKind()); + TEST_CASE(TestImageDataCast()); + + TEST_CASES_END(); +}