diff --git a/include/common/validation_helpers.hpp b/include/common/validation_helpers.hpp index a2e67095..41eef7ad 100644 --- a/include/common/validation_helpers.hpp +++ b/include/common/validation_helpers.hpp @@ -24,6 +24,7 @@ THE SOFTWARE. #include #include +#include "core/image_format.hpp" #include "core/tensor.hpp" /** @@ -83,4 +84,72 @@ THE SOFTWARE. if (std::find(v.begin(), v.end(), tensor.shape(tensor.layout().channels_index())) == v.end()) { \ throw roccv::Exception("Unsupported channel count: " #tensor, eStatusType::INVALID_COMBINATION); \ } \ + } while (0); + +// --------------------------------------------------------------------------- +// ImageBatchVarShape validation +// +// Mirrors the CHECK_TENSOR_* helpers for variable-shape image batches. Datatype +// and channel checks read the batch's uniqueFormat(), which collapses to +// FMT_NONE for an empty or heterogeneous batch — so call +// CHECK_IMAGE_BATCH_UNIFORM_FORMAT first to get a precise error before those. +// --------------------------------------------------------------------------- + +/** + * @brief Validates whether an image batch is located on a specified device. + * + */ +#define CHECK_IMAGE_BATCH_DEVICE(batch, batch_device) \ + if (batch.device() != batch_device) { \ + throw roccv::Exception("Invalid image batch " #batch \ + ": Ensure that this batch is allocated on the proper device.", \ + eStatusType::INVALID_OPERATION); \ + } + +/** + * @brief Validates that an image batch is non-empty. + * + */ +#define CHECK_IMAGE_BATCH_NOT_EMPTY(batch) \ + if (batch.numImages() == 0) { \ + throw roccv::Exception("Image batch " #batch " is empty.", eStatusType::INVALID_COMBINATION); \ + } + +/** + * @brief Validates that every image in the batch shares a single ImageFormat. Fails for an empty or + * heterogeneous batch (uniqueFormat() == FMT_NONE). + * + */ +#define CHECK_IMAGE_BATCH_UNIFORM_FORMAT(batch) \ + if (batch.uniqueFormat() == FMT_NONE) { \ + throw roccv::Exception("Image batch " #batch \ + " must share a single image format across all images (and be non-empty).", \ + eStatusType::INVALID_COMBINATION); \ + } + +/** + * @brief Validates that the batch's shared-format datatype is supported, based on a list of provided + * datatypes. Assumes a uniform format (see CHECK_IMAGE_BATCH_UNIFORM_FORMAT). + * + */ +#define CHECK_IMAGE_BATCH_DATATYPES(batch, ...) \ + do { \ + const std::vector v{__VA_ARGS__}; \ + if (std::find(v.begin(), v.end(), batch.uniqueFormat().dtype()) == v.end()) { \ + throw roccv::Exception("Unsupported data type for image batch: " #batch, eStatusType::NOT_IMPLEMENTED); \ + } \ + } while (0); + +/** + * @brief Validates that the batch's shared-format channel count is supported, based on a list of provided + * channel counts. Assumes a uniform format (see CHECK_IMAGE_BATCH_UNIFORM_FORMAT). + * + */ +#define CHECK_IMAGE_BATCH_CHANNELS(batch, ...) \ + do { \ + const std::vector v{__VA_ARGS__}; \ + if (std::find(v.begin(), v.end(), batch.uniqueFormat().channels()) == v.end()) { \ + throw roccv::Exception("Unsupported channel count for image batch: " #batch, \ + eStatusType::INVALID_COMBINATION); \ + } \ } while (0); \ No newline at end of file diff --git a/include/core/wrappers/border_wrapper.hpp b/include/core/wrappers/border_wrapper.hpp index 6289590b..29e88a42 100644 --- a/include/core/wrappers/border_wrapper.hpp +++ b/include/core/wrappers/border_wrapper.hpp @@ -22,7 +22,7 @@ #pragma once #include "core/detail/sampling_helpers.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "operator_types.h" namespace roccv { @@ -87,37 +87,35 @@ __device__ __host__ inline int64_t reflect101_border_coord_i64(int64_t coord, in } // namespace detail /** - * @brief Wrapper class for ImageWrapper. This extends the descriptors by defining behaviors for when tensor - * coordinates go out of scope. + * @brief Wrapper class which adds border-handling behavior on top of an underlying image wrapper. + * + * Templated on the wrapper type W (e.g. TensorWrapper, ImageBatchVarShapeWrapper) so the same border math + * serves both uniform-shape and variable-shape image batches. The pixel value type T is recovered from + * W::ValueType. W must expose: ValueType, at(n,h,w,c), width(n), height(n), batches(), channels(). * - * @tparam T The underlying data type of the tensor. * @tparam BorderType The border type to use when coordinates are out of bounds. + * @tparam W The underlying image wrapper type. */ -template +template class BorderWrapper { public: - /** - * @brief Wraps an ImageWrapper and extends its capabilities to handle out of bounds coordinates. - * - * @param tensor The tensor to wrap. - * @param border_value The fallback border color to use when using a constant border mode. - */ - BorderWrapper(const Tensor& tensor, T border_value) : m_desc(tensor), m_border_value(border_value) {} + using ValueType = typename W::ValueType; + using WrapperType = W; + static constexpr eBorderType kBorderType = BorderType; /** - * @brief Constructs a BorderWrapper from an existing ImageWrapper. Extends its capabilities to handle out of bound - * coordinates. + * @brief Constructs a BorderWrapper from an existing image wrapper. Extends its capabilities to handle out of + * bound coordinates. * - * @param image_wrapper The ImageWrapper to wrap around the BorderWrapper. - * @param border_value The fallback border color to use when using a constant border mode. + * @param image_wrapper The image wrapper to wrap. + * @param border_value The fallback border color to use when using a constant border mode. */ - BorderWrapper(ImageWrapper image_wrapper, T border_value) - : m_desc(image_wrapper), m_border_value(border_value) {} + BorderWrapper(W image_wrapper, ValueType border_value) : m_desc(image_wrapper), m_border_value(border_value) {} /** * @brief Sample the underlying image with no border logic. Caller must ensure coordinates are in-range. */ - __device__ __host__ inline const T at_inbounds(int64_t n, int64_t h, int64_t w, int64_t c) const { + __device__ __host__ inline const ValueType at_inbounds(int64_t n, int64_t h, int64_t w, int64_t c) const { return m_desc.at(n, h, w, c); } @@ -131,11 +129,14 @@ class BorderWrapper { * @param c The channel index. * @return A reference to the underlying data or a fallback border value of type T. */ - __device__ __host__ const T at(int64_t n, int64_t h, int64_t w, int64_t c) const { + __device__ __host__ const ValueType at(int64_t n, int64_t h, int64_t w, int64_t c) const { + const int64_t imgWidth = width(n); + const int64_t imgHeight = height(n); + // Constant border type implementation. This is a special case which doesn't remap values, but rather returns // the provided constant value. if constexpr (BorderType == eBorderType::BORDER_TYPE_CONSTANT) { - if (w < 0 || w >= width() || h < 0 || h >= height()) + if (w < 0 || w >= imgWidth || h < 0 || h >= imgHeight) return m_border_value; else return m_desc.at(n, h, w, c); @@ -145,13 +146,12 @@ class BorderWrapper { // required at image borders. While this may cause branch divergence, a good bulk of the pixels should fall // within image bounds and will take the same branch. This is preferred over having to do expensive calculations // for EVERY pixel in the image (most of which do not require said calculations). - if (w >= 0 && w < width() && h >= 0 && h < height()) { + if (w >= 0 && w < imgWidth && h >= 0 && h < imgHeight) { return m_desc.at(n, h, w, c); } // Otherwise, do some additional calculations to map the provided x and y coordinates to be within bounds. int64_t x = w, y = h; - int64_t imgWidth = width(), imgHeight = height(); // Reflect border type implementation. (Note: This is NOT REFLECT101, pixels at the border will be duplicated as // is the intended behavior for this border mode.) @@ -185,18 +185,20 @@ class BorderWrapper { } /** - * @brief Retrives the height of the images. + * @brief Retrives the height of the image at batch index n. * + * @param n Batch index. Ignored when W is a uniform-shape wrapper. * @return Image height. */ - __device__ __host__ inline int64_t height() const { return m_desc.height(); } + __device__ __host__ inline int64_t height(int64_t n = 0) const { return m_desc.height(n); } /** - * @brief Retrieves the width of the image. + * @brief Retrieves the width of the image at batch index n. * + * @param n Batch index. Ignored when W is a uniform-shape wrapper. * @return Image width. */ - __device__ __host__ inline int64_t width() const { return m_desc.width(); } + __device__ __host__ inline int64_t width(int64_t n = 0) const { return m_desc.width(n); } /** * @brief Retrieves the number of batches in the image tensor. @@ -213,7 +215,21 @@ class BorderWrapper { __device__ __host__ inline int64_t channels() const { return m_desc.channels(); } private: - ImageWrapper m_desc; - T m_border_value; + W m_desc; + ValueType m_border_value; }; -} // namespace roccv \ No newline at end of file + +/** + * @brief Factory for BorderWrapper. Deduces the underlying wrapper type W from the argument so callers + * only need to spell the border-mode policy explicitly. + * + * @tparam B The border mode to apply. + * @param wrap The underlying image wrapper. + * @param borderValue Fallback value used when B is BORDER_TYPE_CONSTANT. + */ +template +auto MakeBorderWrapper(W wrap, typename W::ValueType borderValue) { + return BorderWrapper(wrap, borderValue); +} + +} // namespace roccv diff --git a/include/core/wrappers/image_batch_var_shape_wrapper.hpp b/include/core/wrappers/image_batch_var_shape_wrapper.hpp new file mode 100644 index 00000000..b32bbdbb --- /dev/null +++ b/include/core/wrappers/image_batch_var_shape_wrapper.hpp @@ -0,0 +1,119 @@ +/* + * 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/detail/type_traits.hpp" +#include "core/image_batch_data.hpp" +#include "core/image_buffer.hpp" + +namespace roccv { + +/** + * @brief ImageBatchVarShapeWrapper is a non-owning, kernel-friendly view over an ImageBatchVarShape's + * descriptor table. It satisfies the same wrapper concept as TensorWrapper + * (ValueType, at(n,h,w,c), width(n), height(n), batches(), channels()) so it composes with + * BorderWrapper / InterpolationWrapper unchanged. + * + * Single-plane interleaved (NHWC-style) only — ImageBatchVarShape rejects multi-plane images + * at pushBack, and channel count is derived from T via detail::NumElements. + * + * Pointer residency follows the snapshot it is built from: for a GPU batch + * (ImageBatchVarShapeDataStridedHip) m_imageList is a device pointer, so the wrapper is usable from + * device code; for a CPU batch (ImageBatchVarShapeDataStridedHost) it is a host pointer, usable from + * host code. Every at()/width()/height() call dereferences it, so the caller must run the wrapper on + * the side matching the snapshot's residency. + * + * @tparam T The datatype of an individual pixel (e.g. uchar1, uchar3, uchar4, float1, float4). + */ +template +class ImageBatchVarShapeWrapper { + public: + using ValueType = T; + using BaseType = detail::BaseType; + + ImageBatchVarShapeWrapper() = default; + + /** + * @brief Creates a ImageBatchVarShapeWrapper from a varshape batch data snapshot. + * + * Accepts either residency through the common ImageBatchVarShapeDataStrided base — a GPU + * (...Hip) snapshot yields a device-usable wrapper, a CPU (...Host) snapshot a host-usable one. + * + * @param data The exported descriptor table from ImageBatchVarShape::exportData(stream). + */ + __host__ ImageBatchVarShapeWrapper(const ImageBatchVarShapeDataStrided& data) + : m_imageList(data.imageList()), m_numImages(data.numImages()) {} + + /** + * @brief Returns a reference to data at given image-batch coordinates. + * + * @param n Batch index. + * @param h Row index within image n. + * @param w Column index within image n. + * @param c Channel index within the pixel. + * @return A reference to the underlying pixel-channel value. + */ + __device__ __host__ T& at(int64_t n, int64_t h, int64_t w, int64_t c) { return *doGetPtr(n, h, w, c); } + + __device__ __host__ const T at(int64_t n, int64_t h, int64_t w, int64_t c) const { return *doGetPtr(n, h, w, c); } + + /** + * @brief Width of the image at batch index n. + */ + __device__ __host__ inline int64_t width(int64_t n) const { return m_imageList[n].planes[0].width; } + + /** + * @brief Height of the image at batch index n. + */ + __device__ __host__ inline int64_t height(int64_t n) const { return m_imageList[n].planes[0].height; } + + /** + * @brief Number of images in the batch. + */ + __device__ __host__ inline int64_t batches() const { return m_numImages; } + + /** + * @brief Number of channels per pixel. Derived from T, identical across all images in v1. + */ + __device__ __host__ inline int64_t channels() const { return detail::NumElements; } + + private: + __device__ __host__ inline T* doGetPtr(int64_t n, int64_t h, int64_t w, int64_t c) const { + // Single-plane interleaved NHWC layout: pixel stride is sizeof(T), channel stride is + // sizeof(BaseType). Match TensorWrapper::at semantics — returns a T* offset to (h, w) + // and additionally shifted by c channels. + const ImagePlaneStrided& p = m_imageList[n].planes[0]; + unsigned char* addr = + reinterpret_cast(p.basePtr) + h * p.rowStride + w * sizeof(T) + c * sizeof(BaseType); + return reinterpret_cast(addr); + } + + const ImageBufferStrided* m_imageList = nullptr; + int32_t m_numImages = 0; +}; + +} // namespace roccv diff --git a/include/core/wrappers/interpolation_wrapper.hpp b/include/core/wrappers/interpolation_wrapper.hpp index 48fed6c3..51559f73 100644 --- a/include/core/wrappers/interpolation_wrapper.hpp +++ b/include/core/wrappers/interpolation_wrapper.hpp @@ -29,34 +29,31 @@ namespace roccv { /** - * @brief A kernel-friendly wrapper which provides interpolation logic based on the given - * coordinates. This tensor wrapper is typically only used for input tensors and does not provide write access to its - * underlying data. + * @brief A kernel-friendly wrapper which provides interpolation logic on top of a BorderWrapper. * - * @tparam T Underlying data type of the tensor data. - * @tparam C Number of channels in data type. - * @tparam B Border type to use for interpolation. - * @tparam I Interpolation type to use. + * Templated directly on the BorderWrapper type so the redundant border-mode and underlying-wrapper template + * parameters need only be spelled once (in the BorderWrapper type). Recover the border mode via + * BW::kBorderType and the underlying wrapper type via BW::WrapperType. + * + * Read-only access; do not use for output tensors. + * + * @tparam I Interpolation type to use. + * @tparam BW The BorderWrapper type to wrap. Must expose ValueType plus at(n,h,w,c), width(n), height(n). */ -template +template class InterpolationWrapper { public: - /** - * @brief Wraps a roccv::Tensor in an InterpolationWrapper to provide pixel interpolation when accessing - * non-integer coordinate mappings. - * - * @param tensor The tensor to wrap. - * @param border_value A fallback border value to use in the case of a constant border mode. - */ - InterpolationWrapper(const Tensor& tensor, T border_value) : m_desc(tensor, border_value) {} + using ValueType = typename BW::ValueType; + using BorderType = BW; + static constexpr eInterpolationType kInterpolationType = I; /** - * @brief Wraps a BorderWrapper in an Interpolation wrapper. Extends capabilities to interpolate pixel values when - * given non-integer coordinates. + * @brief Wraps a BorderWrapper in an InterpolationWrapper. Extends capabilities to interpolate pixel values + * when given non-integer coordinates. * * @param borderWrapper The BorderWrapper to wrap. */ - InterpolationWrapper(BorderWrapper borderWrapper) : m_desc(borderWrapper) {} + InterpolationWrapper(BW borderWrapper) : m_desc(borderWrapper) {} /** * @brief This function calculates the weighting coefficients for the Catmull-Rom cubic interpolation. @@ -81,7 +78,7 @@ class InterpolationWrapper { * @param w Width coordinates. * @return An interpolated value. */ - inline __device__ __host__ const T at(int64_t n, float h, float w, int64_t c) const { + inline __device__ __host__ ValueType at(int64_t n, float h, float w, int64_t c) const { if constexpr (I == eInterpolationType::INTERP_TYPE_NEAREST) { return m_desc.at(n, detail::interp_nearest_i64(h), detail::interp_nearest_i64(w), c); } else if constexpr (I == eInterpolationType::INTERP_TYPE_LINEAR) { @@ -90,7 +87,7 @@ class InterpolationWrapper { // - - // v3 -- v4 - using WorkType = detail::MakeType>; + using WorkType = detail::MakeType>; const int64_t x0 = detail::interp_floor_i64(w); const int64_t y0 = detail::interp_floor_i64(h); @@ -101,7 +98,7 @@ class InterpolationWrapper { const float omfx = 1.f - fx; const float omfy = 1.f - fy; - if (x0 >= 0 && y0 >= 0 && x1 < m_desc.width() && y1 < m_desc.height()) { + if (x0 >= 0 && y0 >= 0 && x1 < m_desc.width(n) && y1 < m_desc.height(n)) { auto v1 = detail::RangeCast(m_desc.at_inbounds(n, y0, x0, c)); auto v2 = detail::RangeCast(m_desc.at_inbounds(n, y0, x1, c)); auto v3 = detail::RangeCast(m_desc.at_inbounds(n, y1, x0, c)); @@ -109,7 +106,7 @@ class InterpolationWrapper { auto q1 = v1 * omfx + v2 * fx; auto q2 = v3 * omfx + v4 * fx; auto q = q1 * omfy + q2 * fy; - return detail::RangeCast(q); + return detail::RangeCast(q); } auto v1 = detail::RangeCast(m_desc.at(n, y0, x0, c)); @@ -121,10 +118,10 @@ class InterpolationWrapper { auto q2 = v3 * omfx + v4 * fx; auto q = q1 * omfy + q2 * fy; - return detail::RangeCast(q); + return detail::RangeCast(q); } else if constexpr (I == eInterpolationType::INTERP_TYPE_CUBIC) { using namespace roccv::detail; - using WorkType = detail::MakeType>; + using WorkType = detail::MakeType>; const int64_t int_x = detail::interp_floor_i64(w); const int64_t int_y = detail::interp_floor_i64(h); @@ -145,7 +142,7 @@ class InterpolationWrapper { WorkType sum = SetAll(0.0f); const bool cubic_fast = - int_x >= 1 && int_y >= 1 && (int_x + 2) < m_desc.width() && (int_y + 2) < m_desc.height(); + int_x >= 1 && int_y >= 1 && (int_x + 2) < m_desc.width(n) && (int_y + 2) < m_desc.height(n); k = 0; if (cubic_fast) { #pragma unroll @@ -168,16 +165,29 @@ class InterpolationWrapper { } } - return detail::RangeCast(sum); + return detail::RangeCast(sum); } } - __device__ __host__ inline int64_t height() const { return m_desc.height(); } - __device__ __host__ inline int64_t width() const { return m_desc.width(); } + __device__ __host__ inline int64_t height(int64_t n = 0) const { return m_desc.height(n); } + __device__ __host__ inline int64_t width(int64_t n = 0) const { return m_desc.width(n); } __device__ __host__ inline int64_t batches() const { return m_desc.batches(); } __device__ __host__ inline int64_t channels() const { return m_desc.channels(); } private: - BorderWrapper m_desc; + BW m_desc; }; + +/** + * @brief Factory for InterpolationWrapper. Deduces the BorderWrapper type BW (and its border mode + + * underlying wrapper) from the argument; callers only need to spell the interpolation policy. + * + * @tparam I The interpolation type. + * @param borderWrap An already-constructed BorderWrapper (typically via MakeBorderWrapper(...)). + */ +template +auto MakeInterpolationWrapper(BW borderWrap) { + return InterpolationWrapper(borderWrap); +} + } // namespace roccv \ No newline at end of file diff --git a/include/core/wrappers/image_wrapper.hpp b/include/core/wrappers/tensor_wrapper.hpp similarity index 81% rename from include/core/wrappers/image_wrapper.hpp rename to include/core/wrappers/tensor_wrapper.hpp index e174c64a..be776576 100644 --- a/include/core/wrappers/image_wrapper.hpp +++ b/include/core/wrappers/tensor_wrapper.hpp @@ -31,28 +31,30 @@ namespace roccv { /** - * @brief ImageWrapper is a non-owning wrapper for roccv::Tensors with a NHWC/NCHW/HWC layout. It provides + * @brief TensorWrapper is a non-owning wrapper for roccv::Tensors with a NHWC/NCHW/HWC layout. It provides * methods for accessing the underlying data within HIP kernels. * * @tparam T The datatype of the underlying tensor data. */ template -class ImageWrapper { +class TensorWrapper { public: using ValueType = T; using BaseType = detail::BaseType; + TensorWrapper() = default; + /** - * @brief Creates an ImageWrapper from a Tensor. + * @brief Creates an TensorWrapper from a Tensor. * - * @param tensor The Tensor to be represented by the ImageWrapper. + * @param tensor The Tensor to be represented by the TensorWrapper. */ - ImageWrapper(const Tensor& tensor) { + TensorWrapper(const Tensor& tensor) { if (tensor.layout() != eTensorLayout::TENSOR_LAYOUT_NHWC && tensor.layout() != eTensorLayout::TENSOR_LAYOUT_NCHW && tensor.layout() != eTensorLayout::TENSOR_LAYOUT_HWC && tensor.layout() != eTensorLayout::TENSOR_LAYOUT_CHW) { - throw Exception("The given tensor layout is not supported for ImageWrapper", eStatusType::NOT_IMPLEMENTED); + throw Exception("The given tensor layout is not supported for TensorWrapper", eStatusType::NOT_IMPLEMENTED); } // Copy tensor data into image tensor descriptor @@ -71,14 +73,14 @@ class ImageWrapper { } /** - * @brief Creates an ImageWrapper from a vector. + * @brief Creates an TensorWrapper from a vector. * * @param input The input vector to wrap. * @param batchSize The number of images within the batch. * @param width The width of each image within the batch. * @param height The height of each image within the batch. */ - ImageWrapper(std::vector& input, int32_t batchSize, int32_t width, int32_t height) { + TensorWrapper(std::vector& input, int32_t batchSize, int32_t width, int32_t height) { // Calculate strides based on input (byte-wise strides) stride.c = sizeof(BaseType); stride.w = stride.c * detail::NumElements; @@ -96,14 +98,14 @@ class ImageWrapper { } /** - * @brief Creates an ImageWrapper from a pointer. + * @brief Creates an TensorWrapper from a pointer. * * @param input The input pointer to wrap. * @param batchSize The number of images within the batch. * @param width The width of each image within the batch. * @param height The height of each image within the batch. */ - ImageWrapper(void* input, int32_t batchSize, int32_t width, int32_t height) { + TensorWrapper(void* input, int32_t batchSize, int32_t width, int32_t height) { // Calculate strides based on input (byte-wise strides) stride.c = sizeof(BaseType); stride.w = stride.c * detail::NumElements; @@ -139,16 +141,22 @@ class ImageWrapper { /** * @brief Retrives the height of the images. * + * @param n Batch index. Ignored for uniform-shape TensorWrapper; included so the signature matches + * ImageBatchVarShapeWrapper, allowing both to satisfy the wrapper concept consumed by + * BorderWrapper / InterpolationWrapper. * @return Image height. */ - __device__ __host__ inline int64_t height() const { return shape.h; } + __device__ __host__ inline int64_t height(int64_t /*n*/ = 0) const { return shape.h; } /** * @brief Retrieves the width of the image. * + * @param n Batch index. Ignored for uniform-shape TensorWrapper; included so the signature matches + * ImageBatchVarShapeWrapper, allowing both to satisfy the wrapper concept consumed by + * BorderWrapper / InterpolationWrapper. * @return Image width. */ - __device__ __host__ inline int64_t width() const { return shape.w; } + __device__ __host__ inline int64_t width(int64_t /*n*/ = 0) const { return shape.w; } /** * @brief Retrieves the number of batches in the image tensor. diff --git a/include/kernels/device/brightness_contrast_device.hpp b/include/kernels/device/brightness_contrast_device.hpp index bccc11d6..b1ac7b4f 100644 --- a/include/kernels/device/brightness_contrast_device.hpp +++ b/include/kernels/device/brightness_contrast_device.hpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" namespace Kernels { namespace Device { diff --git a/include/kernels/device/convert_to_device.hpp b/include/kernels/device/convert_to_device.hpp index 67596f36..4e20be42 100644 --- a/include/kernels/device/convert_to_device.hpp +++ b/include/kernels/device/convert_to_device.hpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" namespace Kernels { namespace Device { diff --git a/include/kernels/device/copy_make_border_device.hpp b/include/kernels/device/copy_make_border_device.hpp index aeae2d38..7fd16461 100644 --- a/include/kernels/device/copy_make_border_device.hpp +++ b/include/kernels/device/copy_make_border_device.hpp @@ -29,9 +29,9 @@ namespace Device { * @brief GPU kernel for CopyMakeBorder operator. * * @tparam SrcDesc Must be a BorderWrapper. - * @tparam DstDesc Must be a ImageWrapper. + * @tparam DstDesc Must be a TensorWrapper. * @param src A BorderWrapper containing information for the input tensor. - * @param dst A ImageWrapper containing information for the output tensor. + * @param dst A TensorWrapper containing information for the output tensor. * @param top The top pixel coordinate on the y-axis where the border should start. * @param left The left-most pixel coordinate on the x-axis where the border should start. * @return __global__ diff --git a/include/kernels/device/reformat_device.hpp b/include/kernels/device/reformat_device.hpp index 35054752..95d6c32a 100644 --- a/include/kernels/device/reformat_device.hpp +++ b/include/kernels/device/reformat_device.hpp @@ -21,7 +21,7 @@ #include -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" namespace Kernels::Device { @@ -34,7 +34,7 @@ namespace Kernels::Device { * @param[out] output The output tensor. */ template -__global__ void reformat(roccv::ImageWrapper input, roccv::ImageWrapper output) { +__global__ void reformat(roccv::TensorWrapper input, roccv::TensorWrapper output) { const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; const int b = blockIdx.z; diff --git a/include/kernels/device/resize_device.hpp b/include/kernels/device/resize_device.hpp index 8bb42033..bbaa4ebe 100644 --- a/include/kernels/device/resize_device.hpp +++ b/include/kernels/device/resize_device.hpp @@ -24,14 +24,20 @@ #include namespace Kernels::Device { +// The scale factor is derived per batch index from input.width(batch)/output.width(): for a uniform +// TensorWrapper input width(n) ignores n and yields the same scale for every image (identical to the prior +// host-precomputed scale), while for an ImageBatchVarShapeWrapper input it yields each image's own scale. template -__global__ void resize(SrcWrapper input, DstWrapper output, float scaleX, float scaleY) { +__global__ void resize(SrcWrapper input, DstWrapper output) { const int x = blockIdx.x * blockDim.x + threadIdx.x; const int y = blockIdx.y * blockDim.y + threadIdx.y; const int batch = blockIdx.z; if (x >= output.width() || y >= output.height()) return; + float scaleX = input.width(batch) / static_cast(output.width()); + float scaleY = input.height(batch) / static_cast(output.height()); + float srcX = fmaf(x + 0.5f, scaleX, -0.5f); float srcY = fmaf(y + 0.5f, scaleY, -0.5f); output.at(batch, y, x, 0) = input.at(batch, srcY, srcX, 0); diff --git a/include/kernels/host/brightness_contrast_host.hpp b/include/kernels/host/brightness_contrast_host.hpp index 16783423..7ee9c40b 100644 --- a/include/kernels/host/brightness_contrast_host.hpp +++ b/include/kernels/host/brightness_contrast_host.hpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" namespace Kernels { namespace Host { diff --git a/include/kernels/host/convert_to_host.hpp b/include/kernels/host/convert_to_host.hpp index 93c19521..13b8381c 100644 --- a/include/kernels/host/convert_to_host.hpp +++ b/include/kernels/host/convert_to_host.hpp @@ -25,7 +25,7 @@ THE SOFTWARE. #include #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" namespace Kernels { namespace Host { diff --git a/include/kernels/host/reformat_host.hpp b/include/kernels/host/reformat_host.hpp index 69e21b67..00980104 100644 --- a/include/kernels/host/reformat_host.hpp +++ b/include/kernels/host/reformat_host.hpp @@ -23,7 +23,7 @@ #include -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" namespace Kernels::Host { @@ -36,7 +36,7 @@ namespace Kernels::Host { * @param[out] output The output tensor. */ template -void reformat(roccv::ImageWrapper input, roccv::ImageWrapper output) { +void reformat(roccv::TensorWrapper input, roccv::TensorWrapper output) { #pragma omp parallel for for (int b = 0; b < output.batches(); b++) { for (int y = 0; y < output.height(); y++) { diff --git a/include/kernels/host/resize_host.hpp b/include/kernels/host/resize_host.hpp index 74aa85b1..463469dc 100644 --- a/include/kernels/host/resize_host.hpp +++ b/include/kernels/host/resize_host.hpp @@ -24,10 +24,14 @@ #include namespace Kernels::Host { +// Scale is derived per batch index (input.width(batch)/output.width()): uniform for a TensorWrapper input, +// per-image for an ImageBatchVarShapeWrapper input. See the device kernel for the rationale. template -void resize(SrcWrapper input, DstWrapper output, float scaleX, float scaleY) { +void resize(SrcWrapper input, DstWrapper output) { #pragma omp parallel for for (int batch = 0; batch < output.batches(); batch++) { + float scaleX = input.width(batch) / static_cast(output.width()); + float scaleY = input.height(batch) / static_cast(output.height()); for (int y = 0; y < output.height(); y++) { for (int x = 0; x < output.width(); x++) { float srcX = fmaf(x + 0.5f, scaleX, -0.5f); diff --git a/include/op_resize.hpp b/include/op_resize.hpp index 868d2238..5bda2cc6 100644 --- a/include/op_resize.hpp +++ b/include/op_resize.hpp @@ -21,6 +21,7 @@ THE SOFTWARE. */ #pragma once +#include "core/image_batch_var_shape.hpp" #include "core/tensor.hpp" #include "i_operator.hpp" #include "operator_types.h" @@ -70,7 +71,7 @@ class Resize final : public IOperator { * Height | No * Batch | Yes * - * Supported interpolation modes: [NEAREST, LINEAR] + * Supported interpolation modes: [NEAREST, LINEAR, CUBIC] * * @param[in] stream The HIP stream to run this operator on. * @param[in] input Input tensor with image batch data. @@ -80,5 +81,46 @@ class Resize final : public IOperator { */ void operator()(hipStream_t stream, const Tensor &in, const Tensor &output, eInterpolationType interpolation, eDeviceType device = eDeviceType::GPU) const; + + /** + * @brief Resizes a batch of variable-sized images into a uniform, constant-sized output tensor. + * + * Every image in the batch is independently resized to the output tensor's per-image width/height using the + * given interpolation mode. The batch must share a single ImageFormat (uniqueFormat() != FMT_NONE); the output + * tensor's dtype, channel count, and batch dimension must match the batch's format and image count. + * + * Limitations: + * + * Input (ImageBatchVarShape): + * Channels: [1, 3, 4] + * Supported DataType(s): [U8, F32] + * Image format must be uniform across the batch. + * + * Output (Tensor): + * Supported TensorLayout(s): [NHWC] + * Channels: [1, 3, 4] + * Supported DataType(s) [U8, F32] + * + * Input/Output dependency: + * + * Property | Input == Output + * -------------- | ------------- + * DataType | Yes + * Channels | Yes + * Width | No + * Height | No + * Batch | Yes (numImages == output N) + * + * Supported interpolation modes: [NEAREST, LINEAR, CUBIC] + * + * @param[in] stream The HIP stream to run this operator on. + * @param[in] input Variable-shape image batch to resize. Non-const because exporting its descriptor snapshot + * advances the batch's lazy device-sync state. + * @param[out] output Constant-sized output tensor (NHWC) receiving the resized batch. + * @param[in] interpolation The interpolation method used when resizing images. + * @param[in] device The device to run this operator on. (Default: GPU). + */ + void operator()(hipStream_t stream, ImageBatchVarShape &input, const Tensor &output, + eInterpolationType interpolation, eDeviceType device = eDeviceType::GPU) const; }; } // namespace roccv diff --git a/python/include/py_image.hpp b/python/include/py_image.hpp new file mode 100644 index 00000000..5c1ebe28 --- /dev/null +++ b/python/include/py_image.hpp @@ -0,0 +1,192 @@ +/** +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 py = pybind11; + +// Python-layer size type: a plain std::tuple rather than a bound +// class. pybind11 auto-converts Python tuples/lists to it, and the binding +// converts to roccv::Size2D internally. +using PySize2D = std::tuple; + +/** + * @brief Python container wrapping a single variable-sized roccv::Image. + * + * Mirrors the role PyTensor plays for roccv::Tensor: a refcounted handle that + * keeps an underlying roccv::Image alive for the Python runtime and exposes its + * metadata plus zero-copy interop. PyImage is the per-element type pushed into a + * PyImageBatchVarShape. + * + * Two construction paths exist: + * - Allocating: build a fresh device buffer from (size, format). + * - Wrapping: adopt an externally-owned buffer exposed via DLPack (zero-copy), + * in which case lifetime is governed by the producer's DLManagedTensor + * deleter rather than by roccv::Image's own allocation. + */ +class PyImage : public std::enable_shared_from_this { + public: + /** + * @brief Allocates a new image buffer of the given size and format on the + * specified device, wrapping the resulting roccv::Image. + * + * @param size The image dimensions (width, height) in pixels. + * @param format The pixel format (dtype + channel count + swizzle). + * @param device The device the image buffer is allocated on. + */ + PyImage(PySize2D size, roccv::ImageFormat format, eDeviceType device); + + /** + * @brief Wraps an existing roccv::Image inside a newly constructed PyImage. + * + * @param image A shared pointer to the roccv::Image to wrap. + */ + PyImage(std::shared_ptr image); + + /** + * @brief Wraps a roccv::Image together with a DLManagedTensor consumed from + * another framework. The underlying roccv::Image must reference the wrapped + * buffer view-only; ownership is released back through the producer's + * DLManagedTensor.deleter, invoked by this object's destructor. + * + * @param image A shared pointer to the roccv::Image to wrap. + * @param managedTensor The DLManagedTensor provided by the producer. + */ + PyImage(std::shared_ptr image, DLManagedTensor* managedTensor); + + /** + * @brief Destroys the PyImage and invokes the wrapped DLManagedTensor's + * deleter, if one exists and is valid. + */ + ~PyImage(); + + /** + * @brief Creates a new image of the given size and format filled with zeros. + * + * @param size The image dimensions (width, height) in pixels. + * @param format The pixel format. + * @param device The device the image buffer is allocated on. + * @return std::shared_ptr + */ + static std::shared_ptr Zeros(PySize2D size, roccv::ImageFormat format, eDeviceType device); + + /** + * @brief Wraps an external buffer (a capsule containing a DLManagedTensor) + * as an image without copying. + * + * @param src A capsule containing a DLManagedTensor. + * @param format The pixel format to interpret the buffer as. + * @return std::shared_ptr + */ + static std::shared_ptr WrapExternalBuffer(py::object src, roccv::ImageFormat format); + + /** + * @brief Exports this image as a capsule containing a DLManagedTensor for + * consumption by another framework. + * + * @param stream Optional stream pointer value (used for framework interop). + * @return py::capsule + */ + py::capsule toDLPack(py::object stream); + + /** + * @brief Returns a tuple describing the DLPack device of this image. Index 0 + * is the device type and index 1 is the device id. + * + * @return py::tuple + */ + py::tuple getDLDevice(); + + /** + * @brief Creates a copy of this image on the specified device. + * + * @param device The device of the new image. + * @return std::shared_ptr + */ + std::shared_ptr copyTo(eDeviceType device); + + /** + * @brief Gets the size (width, height) of this image. + * + * @return PySize2D + */ + PySize2D getSize(); + + /** + * @brief Gets the width of this image in pixels. + * + * @return int + */ + int getWidth(); + + /** + * @brief Gets the height of this image in pixels. + * + * @return int + */ + int getHeight(); + + /** + * @brief Gets the pixel format of this image. + * + * @return roccv::ImageFormat + */ + roccv::ImageFormat getFormat(); + + /** + * @brief Gets the device this image resides on. + * + * @return eDeviceType + */ + eDeviceType getDevice(); + + /** + * @brief Gets the underlying roccv::Image this container wraps. + * + * @return std::shared_ptr + */ + std::shared_ptr getImage(); + + /** + * @brief Exports this class in the provided module. + * + * @param m The python module to export this class to. + */ + static void Export(py::module& m); + + private: + std::shared_ptr m_image; + + // TODO: DLManagedTensor is slated for deprecation in future DLPack versions. + // Update this to the versioned managed tensor before then. + DLManagedTensor* m_managedTensor = nullptr; +}; diff --git a/python/include/py_image_batch_var_shape.hpp b/python/include/py_image_batch_var_shape.hpp new file mode 100644 index 00000000..fbd69236 --- /dev/null +++ b/python/include/py_image_batch_var_shape.hpp @@ -0,0 +1,189 @@ +/** +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 "py_image.hpp" + +namespace py = pybind11; + +/** + * @brief Python container wrapping a roccv::ImageBatchVarShape. + * + * Holds a batch of variable-sized images that may differ in size, format, and + * dtype. Capacity is fixed at construction; images are appended via pushBack up + * to that capacity. + * + * The underlying roccv::ImageBatchVarShape is move-only and stores its images by + * value, so this wrapper keeps a parallel vector of PyImage handles. That mirror + * serves two purposes: it keeps each image's Python object (and any wrapped + * DLManagedTensor) alive for as long as the batch references it, and it backs + * Python-level indexing/iteration without rebuilding handles from the underlying + * batch. + */ +class PyImageBatchVarShape : public std::enable_shared_from_this { + public: + /** + * @brief Constructs an empty batch with the given capacity on the specified + * device, wrapping a newly constructed roccv::ImageBatchVarShape. + * + * @param capacity The maximum number of images the batch can hold. + * @param device The device the batch (and every image it accepts) resides on. + */ + PyImageBatchVarShape(int32_t capacity, eDeviceType device); + + ~PyImageBatchVarShape() = default; + + /** + * @brief Creates a new batch with the given capacity on the specified device. + * + * @param capacity The maximum number of images the batch can hold. + * @param device The device the batch resides on. + * @return std::shared_ptr + */ + static std::shared_ptr Create(int32_t capacity, eDeviceType device); + + /** + * @brief Builds a batch by wrapping a vector of external buffers (capsules + * containing DLManagedTensors) as images without copying. Capacity is set to + * the number of buffers provided. + * + * @param buffers A list of capsules, each containing a DLManagedTensor. + * @param format The pixel format to interpret each buffer as. + * @return std::shared_ptr + */ + static std::shared_ptr WrapExternalBufferVector(std::vector buffers, + roccv::ImageFormat format); + + /** + * @brief Appends a single image to the end of the batch. Throws if capacity + * would be exceeded or the image's device does not match the batch's. + * + * @param image The image to append. + */ + void pushBack(std::shared_ptr image); + + /** + * @brief Appends multiple images to the end of the batch. Provides a strong + * exception guarantee: on failure the batch is rolled back to its pre-call + * state. + * + * @param images The images to append. + */ + void pushBackMany(const std::vector>& images); + + /** + * @brief Removes the trailing `count` images from the batch. + * + * @param count The number of images to remove. + */ + void popBack(int32_t count); + + /** + * @brief Removes all images from the batch. Capacity is retained and the + * batch remains reusable. + */ + void clear(); + + /** + * @brief Gets the number of images currently in the batch (__len__). + * + * @return int32_t + */ + int32_t numImages(); + + /** + * @brief Gets the maximum number of images the batch can hold. + * + * @return int32_t + */ + int32_t capacity(); + + /** + * @brief Gets the bounding box (width, height) across all images in pixels. + * + * @return PySize2D + */ + PySize2D maxSize(); + + /** + * @brief Gets the common format across all images, or FMT_NONE if the formats + * are heterogeneous or the batch is empty. + * + * @return roccv::ImageFormat + */ + roccv::ImageFormat uniqueFormat(); + + /** + * @brief Gets the device this batch resides on. + * + * @return eDeviceType + */ + eDeviceType getDevice(); + + /** + * @brief Gets the image at the given index (__getitem__). + * + * @param index The index of the image to retrieve. + * @return std::shared_ptr + */ + std::shared_ptr at(int32_t index); + + /** + * @brief Gets the list of PyImage handles currently held by the batch. Backs + * Python-level iteration over the batch. + * + * @return const std::vector>& + */ + const std::vector>& images() const; + + /** + * @brief Gets the underlying roccv::ImageBatchVarShape this container wraps. + * + * @return std::shared_ptr + */ + std::shared_ptr getBatch(); + + /** + * @brief Exports this class in the provided module. + * + * @param m The python module to export this class to. + */ + static void Export(py::module& m); + + private: + std::shared_ptr m_batch; + + // Keep-alive mirror of the images pushed into m_batch. Keeps each PyImage + // (and any wrapped DLManagedTensor) alive for the batch's lifetime and backs + // Python indexing/iteration. + std::vector> m_images; +}; diff --git a/python/include/py_image_format.hpp b/python/include/py_image_format.hpp new file mode 100644 index 00000000..e4ab54ff --- /dev/null +++ b/python/include/py_image_format.hpp @@ -0,0 +1,41 @@ +/** +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 + +namespace py = pybind11; + +/** + * @brief Exports roccv::ImageFormat to Python as the `Format` type. + * + * Every FMT_* constant is exposed as a named attribute (e.g. Format.RGB8), with + * read-only `channels`/`dtype`/`swizzle`/`planes` properties and a named + * __repr__. Because roccv::ImageFormat is a value class (dtype + channels + + * swizzle) rather than a packed integer, it is bound as a py::class_ instead of + * a py::enum_. + */ +class PyImageFormat { + public: + static void Export(py::module& m); +}; diff --git a/python/src/main.cpp b/python/src/main.cpp index 1bc89b22..5a2195b5 100644 --- a/python/src/main.cpp +++ b/python/src/main.cpp @@ -47,6 +47,9 @@ THE SOFTWARE. #include "operators/py_op_warp_perspective.hpp" #include "py_enums.hpp" #include "py_exception.hpp" +#include "py_image.hpp" +#include "py_image_batch_var_shape.hpp" +#include "py_image_format.hpp" #include "py_stream.hpp" #include "py_structs.hpp" #include "py_tensor.hpp" @@ -60,8 +63,11 @@ PYBIND11_MODULE(rocpycv, m) { PyException::Export(m); PyEnums::Export(m); PyStructs::Export(m); + PyImageFormat::Export(m); PyStream::Export(m); PyTensor::Export(m); + PyImage::Export(m); + PyImageBatchVarShape::Export(m); PyOpCustomCrop::Export(m); PyOpNonMaxSuppression::Export(m); PyOpNormalize::Export(m); diff --git a/python/src/py_enums.cpp b/python/src/py_enums.cpp index 0ce39d08..778b718c 100644 --- a/python/src/py_enums.cpp +++ b/python/src/py_enums.cpp @@ -25,7 +25,15 @@ THE SOFTWARE. #include #include +#include + void PyEnums::Export(py::module& m) { + // eSwizzle Enum Bindings + py::enum_(m, "eSwizzle") + .value("XYZW", roccv::eSwizzle::XYZW) + .value("ZYXW", roccv::eSwizzle::ZYXW) + .export_values(); + // eTensorLayout Enum Bindings py::enum_(m, "eTensorLayout") .value("NHWC", TENSOR_LAYOUT_NHWC) diff --git a/python/src/py_image.cpp b/python/src/py_image.cpp new file mode 100644 index 00000000..37bff753 --- /dev/null +++ b/python/src/py_image.cpp @@ -0,0 +1,269 @@ +/** +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 "py_image.hpp" + +#include +#include + +#include +#include +#include +#include +#include +#include + +#include "py_helpers.hpp" + +namespace { + +// Maps a (source, destination) device pair to the appropriate HIP memcpy kind. +const std::map, hipMemcpyKind> kDevicePairToMemcpyKind = { + {{eDeviceType::CPU, eDeviceType::CPU}, hipMemcpyHostToHost}, + {{eDeviceType::CPU, eDeviceType::GPU}, hipMemcpyHostToDevice}, + {{eDeviceType::GPU, eDeviceType::CPU}, hipMemcpyDeviceToHost}, + {{eDeviceType::GPU, eDeviceType::GPU}, hipMemcpyDeviceToDevice}}; + +// Builds a DLManagedTensor describing an image as an interleaved HWC buffer: +// shape [height, width, channels], element-wise strides, single per-channel +// dtype. The owning PyImage is stashed in manager_ctx so the underlying buffer +// outlives every DLPack consumer. +DLManagedTensor* createDLManagedTensorFromImage(std::shared_ptr image, std::shared_ptr ctx) { + auto imageData = image->exportData(); + const roccv::ImagePlaneStrided& plane = imageData.plane(0); + + const roccv::ImageFormat format = image->format(); + const int64_t channels = format.channels(); + const size_t elemSize = roccv::DataType(format.dtype()).size(); + + DLManagedTensor* dlTensor = new DLManagedTensor(); + dlTensor->dl_tensor.data = plane.basePtr; + dlTensor->dl_tensor.byte_offset = 0; + dlTensor->dl_tensor.ndim = 3; + dlTensor->dl_tensor.device = RoccvDeviceToDLDevice(image->device()); + dlTensor->dl_tensor.dtype = RoccvTypeToDLType(format.dtype()); + + // HWC layout. DLPack strides are element-wise, so the row stride (bytes) is + // converted to elements via the per-channel element size. + dlTensor->dl_tensor.shape = new int64_t[3]{plane.height, plane.width, channels}; + dlTensor->dl_tensor.strides = new int64_t[3]{plane.rowStride / static_cast(elemSize), channels, 1}; + + dlTensor->manager_ctx = new std::shared_ptr(ctx); + dlTensor->deleter = [](DLManagedTensor* mt) { + delete[] mt->dl_tensor.shape; + delete[] mt->dl_tensor.strides; + delete static_cast*>(mt->manager_ctx); + delete mt; + }; + + return dlTensor; +} + +} // namespace + +PyImage::PyImage(PySize2D size, roccv::ImageFormat format, eDeviceType device) { + auto [width, height] = size; + m_image = std::make_shared(roccv::Size2D{width, height}, format, device); +} + +PyImage::PyImage(std::shared_ptr image) : m_image(image) {} + +PyImage::PyImage(std::shared_ptr image, DLManagedTensor* managedTensor) + : m_image(image), m_managedTensor(managedTensor) {} + +PyImage::~PyImage() { + // If we are a consumer of a DLManagedTensor, ensure that we call its deleter. + // The wrapped roccv::Image is view-only, so the buffer is released here. + if (m_managedTensor && m_managedTensor->deleter) { + m_managedTensor->deleter(m_managedTensor); + } +} + +std::shared_ptr PyImage::Zeros(PySize2D size, roccv::ImageFormat format, eDeviceType device) { + auto [width, height] = size; + auto image = std::make_shared(roccv::Size2D{width, height}, format, device); + + auto imageData = image->exportData(); + const roccv::ImagePlaneStrided& plane = imageData.plane(0); + const size_t bytes = static_cast(plane.rowStride) * plane.height; + + if (device == eDeviceType::GPU) { + HIP_VALIDATE_NO_ERRORS(hipMemset(plane.basePtr, 0, bytes)); + } else { + std::memset(plane.basePtr, 0, bytes); + } + + return std::make_shared(image); +} + +std::shared_ptr PyImage::WrapExternalBuffer(py::object src, roccv::ImageFormat format) { + if (!py::hasattr(src, "__dlpack__")) { + throw std::runtime_error("Provided object does not support the DLPack protocol."); + } + + py::capsule dlpackCapsule = src.attr("__dlpack__")(); + if (!PyCapsule_IsValid(dlpackCapsule.ptr(), "dltensor")) { + throw std::runtime_error("Invalid DLPack capsule."); + } + DLManagedTensor* dlManagedTensor = static_cast(dlpackCapsule.get_pointer()); + DLTensor dlTensor = dlManagedTensor->dl_tensor; + + // Mark the capsule as consumed so its deleter will not free the underlying + // data; ownership now flows through the returned PyImage's destructor. + dlpackCapsule.set_name("used_dltensor"); + + // Interpret the buffer as HWC (ndim 3) or HW grayscale (ndim 2). + int64_t height, width, channels; + if (dlTensor.ndim == 3) { + height = dlTensor.shape[0]; + width = dlTensor.shape[1]; + channels = dlTensor.shape[2]; + } else if (dlTensor.ndim == 2) { + height = dlTensor.shape[0]; + width = dlTensor.shape[1]; + channels = 1; + } else { + throw std::runtime_error("DLPack buffer must be 2-D (HW) or 3-D (HWC) to be wrapped as an image."); + } + + // Infer the format from the buffer when one was not supplied. + if (format == roccv::FMT_NONE) { + format = roccv::ImageFormat(DLTypeToRoccvType(dlTensor.dtype), static_cast(channels)); + } + + const size_t elemSize = roccv::DataType(format.dtype()).size(); + + // DLPack strides are element-wise; recover the byte row stride. A null stride + // array denotes a compact (C-contiguous) buffer. + int64_t rowStride; + if (dlTensor.strides != nullptr) { + rowStride = dlTensor.strides[0] * static_cast(elemSize); + } else { + rowStride = width * channels * static_cast(elemSize); + } + + roccv::ImageBufferStrided strided{}; + strided.numPlanes = 1; + strided.planes[0].width = static_cast(width); + strided.planes[0].height = static_cast(height); + strided.planes[0].rowStride = rowStride; + strided.planes[0].basePtr = static_cast(dlTensor.data) + dlTensor.byte_offset; + + eDeviceType device = DLDeviceToRoccvDevice(dlTensor.device); + + // View-only wrap (no cleanup callback): the wrapped buffer is freed by this + // PyImage's destructor via the consumed DLManagedTensor, mirroring PyTensor. + roccv::Image image = (device == eDeviceType::GPU) + ? roccv::ImageWrapData(roccv::ImageDataStridedHip(format, strided)) + : roccv::ImageWrapData(roccv::ImageDataStridedHost(format, strided)); + + return std::make_shared(std::make_shared(image), dlManagedTensor); +} + +py::capsule PyImage::toDLPack(py::object /* stream */) { + // Stream parameter is intentionally unused to support framework DLPack device conversions. + DLManagedTensor* dlTensor = createDLManagedTensorFromImage(m_image, shared_from_this()); + + py::capsule capsule(dlTensor, "dltensor", [](PyObject* self) { + if (PyCapsule_IsValid(self, "used_dltensor")) { + return; // Do nothing if the capsule has been consumed. + } + + DLManagedTensor* managed = static_cast(PyCapsule_GetPointer(self, "dltensor")); + if (managed == nullptr) { + PyErr_WriteUnraisable(self); + return; + } + + if (managed->deleter) { + managed->deleter(managed); + } + }); + + return capsule; +} + +py::tuple PyImage::getDLDevice() { + DLDevice device = RoccvDeviceToDLDevice(m_image->device()); + return py::make_tuple(py::int_(static_cast(device.device_type)), py::int_(static_cast(device.device_id))); +} + +std::shared_ptr PyImage::copyTo(eDeviceType device) { + auto dstImage = std::make_shared(m_image->size(), m_image->format(), device); + + auto srcData = m_image->exportData(); + auto dstData = dstImage->exportData(); + const roccv::ImagePlaneStrided& srcPlane = srcData.plane(0); + const roccv::ImagePlaneStrided& dstPlane = dstData.plane(0); + + const roccv::ImageFormat format = m_image->format(); + const size_t elemSize = roccv::DataType(format.dtype()).size(); + const size_t widthBytes = static_cast(srcPlane.width) * format.channels() * elemSize; + + // A pitched copy correctly handles differing source/destination row strides + // (e.g. a wrapped buffer with arbitrary padding copied into a fresh image). + HIP_VALIDATE_NO_ERRORS(hipMemcpy2D(dstPlane.basePtr, dstPlane.rowStride, srcPlane.basePtr, srcPlane.rowStride, + widthBytes, srcPlane.height, + kDevicePairToMemcpyKind.at({m_image->device(), device}))); + + return std::make_shared(dstImage); +} + +PySize2D PyImage::getSize() { + roccv::Size2D size = m_image->size(); + return std::make_tuple(size.w, size.h); +} + +int PyImage::getWidth() { return m_image->size().w; } + +int PyImage::getHeight() { return m_image->size().h; } + +roccv::ImageFormat PyImage::getFormat() { return m_image->format(); } + +eDeviceType PyImage::getDevice() { return m_image->device(); } + +std::shared_ptr PyImage::getImage() { return m_image; } + +void PyImage::Export(py::module& m) { + using namespace py::literals; + + py::class_>(m, "Image", + "A single variable-sized image with device-resident data.") + .def(py::init(), "size"_a, "format"_a, "device"_a = eDeviceType::GPU, + "Allocate a new image of the given size (width, height) and format.") + .def_static("zeros", &PyImage::Zeros, "size"_a, "format"_a, "device"_a = eDeviceType::GPU, + "Create an image of the given size and format filled with zeros.") + .def("copy_to", &PyImage::copyTo, "device"_a, + "Returns a copy of the image with data copied to the given device.") + .def("__dlpack__", &PyImage::toDLPack, "stream"_a = py::none(), "Creates a DLPack capsule from this image.") + .def("__dlpack_device__", &PyImage::getDLDevice, + "Returns a tuple containing the DLPack device type and id for this image.") + .def_property_readonly("size", &PyImage::getSize, + "Read-only property returning the (width, height) of the image.") + .def_property_readonly("width", &PyImage::getWidth, "Read-only property returning the width of the image.") + .def_property_readonly("height", &PyImage::getHeight, "Read-only property returning the height of the image.") + .def_property_readonly("format", &PyImage::getFormat, "Read-only property returning the format of the image.") + .def_property_readonly("device", &PyImage::getDevice, "Read-only property returning the device of the image."); + + m.def("as_image", &PyImage::WrapExternalBuffer, "buffer"_a, "format"_a = roccv::FMT_NONE, + "Wraps a DLPack-supported buffer as a rocpycv Image without copying."); +} diff --git a/python/src/py_image_batch_var_shape.cpp b/python/src/py_image_batch_var_shape.cpp new file mode 100644 index 00000000..3f2169f7 --- /dev/null +++ b/python/src/py_image_batch_var_shape.cpp @@ -0,0 +1,151 @@ +/** +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 "py_image_batch_var_shape.hpp" + +#include + +#include + +PyImageBatchVarShape::PyImageBatchVarShape(int32_t capacity, eDeviceType device) { + m_batch = std::make_shared(capacity, device); +} + +std::shared_ptr PyImageBatchVarShape::Create(int32_t capacity, eDeviceType device) { + return std::make_shared(capacity, device); +} + +std::shared_ptr PyImageBatchVarShape::WrapExternalBufferVector(std::vector buffers, + roccv::ImageFormat format) { + if (buffers.empty()) { + throw py::value_error("as_images requires a non-empty list of buffers."); + } + + // Wrap each buffer first so the batch's device can be inferred from the + // resulting images (they must all share a device). + std::vector> images; + images.reserve(buffers.size()); + for (auto& buffer : buffers) { + images.push_back(PyImage::WrapExternalBuffer(buffer, format)); + } + + eDeviceType device = images.front()->getDevice(); + auto batch = std::make_shared(static_cast(images.size()), device); + batch->pushBackMany(images); + return batch; +} + +void PyImageBatchVarShape::pushBack(std::shared_ptr image) { + // Mutate the underlying batch first; if it throws (capacity/device mismatch) + // the keep-alive mirror is left untouched. + m_batch->pushBack(*image->getImage()); + m_images.push_back(image); +} + +void PyImageBatchVarShape::pushBackMany(const std::vector>& images) { + const size_t oldSize = m_images.size(); + try { + for (const auto& image : images) { + m_batch->pushBack(*image->getImage()); + m_images.push_back(image); + } + } catch (...) { + // Roll back to the pre-call state for the strong exception guarantee. + const int32_t added = static_cast(m_images.size() - oldSize); + if (added > 0) { + m_batch->popBack(added); + m_images.resize(oldSize); + } + throw; + } +} + +void PyImageBatchVarShape::popBack(int32_t count) { + // popBack validates count before mutating, so the mirror stays consistent. + m_batch->popBack(count); + m_images.erase(m_images.end() - count, m_images.end()); +} + +void PyImageBatchVarShape::clear() { + m_batch->clear(); + m_images.clear(); +} + +int32_t PyImageBatchVarShape::numImages() { return m_batch->numImages(); } + +int32_t PyImageBatchVarShape::capacity() { return m_batch->capacity(); } + +PySize2D PyImageBatchVarShape::maxSize() { + roccv::Size2D size = m_batch->maxSize(); + return std::make_tuple(size.w, size.h); +} + +roccv::ImageFormat PyImageBatchVarShape::uniqueFormat() { return m_batch->uniqueFormat(); } + +eDeviceType PyImageBatchVarShape::getDevice() { return m_batch->device(); } + +std::shared_ptr PyImageBatchVarShape::at(int32_t index) { + // Support Python-style negative indexing. + if (index < 0) { + index += numImages(); + } + if (index < 0 || index >= numImages()) { + throw py::index_error("ImageBatchVarShape index out of range."); + } + return m_images[index]; +} + +const std::vector>& PyImageBatchVarShape::images() const { return m_images; } + +std::shared_ptr PyImageBatchVarShape::getBatch() { return m_batch; } + +void PyImageBatchVarShape::Export(py::module& m) { + using namespace py::literals; + + py::class_>( + m, "ImageBatchVarShape", "A batch of variable-sized images that may differ in size, format, and dtype.") + .def(py::init(), "capacity"_a, "device"_a = eDeviceType::GPU, + "Create an empty batch with the given capacity (maximum number of images).") + .def("pushback", &PyImageBatchVarShape::pushBack, "image"_a, "Append a single image to the end of the batch.") + .def("pushback", &PyImageBatchVarShape::pushBackMany, "images"_a, + "Append multiple images to the end of the batch.") + .def("popback", &PyImageBatchVarShape::popBack, "count"_a = 1, + "Remove one or more images from the end of the batch.") + .def("clear", &PyImageBatchVarShape::clear, "Remove all images from the batch; capacity is retained.") + .def("__len__", &PyImageBatchVarShape::numImages, "Return the number of images currently in the batch.") + .def("__getitem__", &PyImageBatchVarShape::at, "index"_a, "Return the image at the given index.") + .def( + "__iter__", + [](PyImageBatchVarShape& self) { return py::make_iterator(self.images().begin(), self.images().end()); }, + py::keep_alive<0, 1>(), "Return an iterator over the images in the batch.") + .def_property_readonly("capacity", &PyImageBatchVarShape::capacity, + "Read-only property returning the maximum number of images the batch can hold.") + .def_property_readonly("maxsize", &PyImageBatchVarShape::maxSize, + "Read-only property returning the (width, height) bounding box across all images.") + .def_property_readonly("uniqueformat", &PyImageBatchVarShape::uniqueFormat, + "Read-only property returning the common format, or Format.NONE if heterogeneous/empty.") + .def_property_readonly("device", &PyImageBatchVarShape::getDevice, + "Read-only property returning the device the batch resides on."); + + m.def("as_images", &PyImageBatchVarShape::WrapExternalBufferVector, "buffers"_a, "format"_a = roccv::FMT_NONE, + "Wraps a list of DLPack-supported buffers as a rocpycv ImageBatchVarShape without copying."); +} diff --git a/python/src/py_image_format.cpp b/python/src/py_image_format.cpp new file mode 100644 index 00000000..6c91062a --- /dev/null +++ b/python/src/py_image_format.cpp @@ -0,0 +1,144 @@ +/** +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 "py_image_format.hpp" + +#include +#include +#include +#include +#include + +// X-macro listing every named ImageFormat constant. Each entry is +// (python_name, FMT_ suffix). This single list drives both the attribute +// definitions and the __repr__ reverse-lookup. Formats whose name begins with a +// digit get an underscore-prefixed python name (e.g. _2F32) to remain valid +// identifiers. +#define ROCCV_FORMAT_LIST(ENTRY) \ + ENTRY("NONE", NONE) \ + ENTRY("U8", U8) \ + ENTRY("S8", S8) \ + ENTRY("U16", U16) \ + ENTRY("S16", S16) \ + ENTRY("U32", U32) \ + ENTRY("S32", S32) \ + ENTRY("F32", F32) \ + ENTRY("_2F32", 2F32) \ + ENTRY("F64", F64) \ + ENTRY("RGB8", RGB8) \ + ENTRY("BGR8", BGR8) \ + ENTRY("RGBA8", RGBA8) \ + ENTRY("BGRA8", BGRA8) \ + ENTRY("RGBs8", RGBs8) \ + ENTRY("RGBAs8", RGBAs8) \ + ENTRY("RGB16", RGB16) \ + ENTRY("RGBA16", RGBA16) \ + ENTRY("RGBs16", RGBs16) \ + ENTRY("RGBAs16", RGBAs16) \ + ENTRY("RGB32", RGB32) \ + ENTRY("RGBA32", RGBA32) \ + ENTRY("RGBs32", RGBs32) \ + ENTRY("RGBAs32", RGBAs32) \ + ENTRY("RGBf32", RGBf32) \ + ENTRY("RGBAf32", RGBAf32) \ + ENTRY("RGBf64", RGBf64) \ + ENTRY("RGBAf64", RGBAf64) + +namespace { + +// Names of the class-level format constants (Format.RGB8, Format.U8, ...). Used +// to hide them from instance attribute lookup so chains like Format.RGB8.U8 are +// rejected (see __getattribute__ in Export). +const std::unordered_set kConstantNames = { +#define ENTRY(pyname, suffix) pyname, + ROCCV_FORMAT_LIST(ENTRY) +#undef ENTRY +}; + +// Produces a human-readable name for a format, preferring the named constant it +// matches (e.g. "rocpycv.Format.RGB8"). Falls back to a structural description +// for formats constructed directly from (dtype, channels, swizzle). +std::string ImageFormatToString(const roccv::ImageFormat& fmt) { +#define ENTRY(pyname, suffix) \ + if (fmt == roccv::FMT_##suffix) return std::string("rocpycv.Format.") + (pyname); + ROCCV_FORMAT_LIST(ENTRY) +#undef ENTRY + + return "rocpycv.Format(dtype=" + std::to_string(static_cast(fmt.dtype())) + + ", channels=" + std::to_string(fmt.channels()) + + ", swizzle=" + std::to_string(static_cast(fmt.swizzle())) + ")"; +} + +size_t ImageFormatHash(const roccv::ImageFormat& fmt) { + size_t h = std::hash()(static_cast(fmt.dtype())); + h = h * 31 + std::hash()(fmt.channels()); + h = h * 31 + std::hash()(static_cast(fmt.swizzle())); + return h; +} + +} // namespace + +void PyImageFormat::Export(py::module& m) { + using namespace py::literals; + using roccv::ImageFormat; + + py::class_ fmt(m, "Format", "Describes how image pixel data is laid out in memory."); + + fmt.def(py::init(), "dtype"_a, "channels"_a, + "swizzle"_a = roccv::eSwizzle::XYZW, "Construct a format from a data type, channel count, and swizzle.") + .def_property_readonly("channels", &ImageFormat::channels, + "Read-only property that returns the number of color channels in the image.") + .def_property_readonly("dtype", &ImageFormat::dtype, + "Read-only property that returns the data type of each channel.") + .def_property_readonly("swizzle", &ImageFormat::swizzle, + "Read-only property that returns the channel swizzle of the format.") + .def_property_readonly( + "planes", [](const ImageFormat&) { return 1; }, + "Read-only property that returns the number of planes in the image (always 1; rocCV is single-plane).") + .def("__eq__", &ImageFormat::operator==, py::is_operator()) + .def("__ne__", &ImageFormat::operator!=, py::is_operator()) + .def("__hash__", &ImageFormatHash) + .def("__repr__", &ImageFormatToString); + + // Attach each named format constant as a class attribute, e.g. Format.RGB8. +#define ENTRY(pyname, suffix) fmt.attr(pyname) = roccv::FMT_##suffix; + ROCCV_FORMAT_LIST(ENTRY) +#undef ENTRY + + // The constants above are themselves Format instances and live in the class + // dict, so ordinary instance attribute lookup would re-resolve them through + // any Format value — enabling nonsensical chains like Format.RGB8.RGB8.U8. + // Override __getattribute__ to hide the constant names from instance lookup; + // class-level access (Format.RGB8) goes through the metaclass and is + // unaffected. This mirrors the instance-isolation that py::enum_ members get. + fmt.def("__getattribute__", [](py::handle self, py::str name) -> py::object { + if (kConstantNames.count(static_cast(name))) { + throw py::attribute_error("'rocpycv.Format' object has no attribute '" + static_cast(name) + + "'"); + } + PyObject* result = PyObject_GenericGetAttr(self.ptr(), name.ptr()); + if (!result) throw py::error_already_set(); + return py::reinterpret_steal(result); + }); +} + +#undef ROCCV_FORMAT_LIST diff --git a/python/src/rocpycv.pyi b/python/src/rocpycv.pyi index 855a8d5a..eae37b18 100644 --- a/python/src/rocpycv.pyi +++ b/python/src/rocpycv.pyi @@ -8,7 +8,7 @@ from __future__ import annotations import collections.abc import typing -__all__: list[str] = ['BGR', 'BINARY', 'BINARY_INV', 'BOTH', 'BT2020', 'BT601', 'BT709', 'BndBox', 'BndBoxes', 'Box', 'CHW', 'COLOR_BGR2GRAY', 'COLOR_BGR2RGB', 'COLOR_BGR2YUV', 'COLOR_BGR2YUV_NV12', 'COLOR_BGR2YUV_NV21', 'COLOR_RGB2BGR', 'COLOR_RGB2GRAY', 'COLOR_RGB2YUV', 'COLOR_RGB2YUV_NV12', 'COLOR_RGB2YUV_NV21', 'COLOR_YUV2BGR', 'COLOR_YUV2BGR_NV12', 'COLOR_YUV2BGR_NV21', 'COLOR_YUV2RGB', 'COLOR_YUV2RGB_NV12', 'COLOR_YUV2RGB_NV21', 'CONSTANT', 'CPU', 'CUBIC', 'ColorRGBA', 'Exception', 'F32', 'F64', 'GPU', 'Grayscale', 'HWC', 'LINEAR', 'N', 'NC', 'NCHW', 'NEAREST', 'NHWC', 'NW', 'NWC', 'NormalizeFlags', 'REFLECT', 'REFLECT101', 'REMAP_ABSOLUTE', 'REMAP_ABSOLUTE_NORMALIZED', 'REMAP_RELATIVE_NORMALIZED', 'REPLICATE', 'RGB', 'S16', 'S32', 'S8', 'Size2D', 'Stream', 'TOZERO', 'TOZERO_INV', 'TRUNC', 'Tensor', 'U16', 'U32', 'U8', 'WRAP', 'X', 'Y', 'YUV', 'YVU', 'advcvtcolor', 'advcvtcolor_into', 'bilateral_filter', 'bilateral_filter_into', 'bndbox', 'bndbox_into', 'center_crop', 'center_crop_into', 'composite', 'composite_into', 'convert_to', 'convert_to_into', 'copymakeborder', 'copymakeborder_into', 'custom_crop', 'custom_crop_into', 'cvtcolor', 'cvtcolor_into', 'eAxis', 'eBorderType', 'eChannelType', 'eColorConversionCode', 'eColorSpec', 'eDataType', 'eDeviceType', 'eInterpolationType', 'eRemapType', 'eTensorLayout', 'eThresholdType', 'flip', 'flip_into', 'from_dlpack', 'gamma_contrast', 'gamma_contrast_into', 'histogram', 'histogram_into', 'nms', 'nms_into', 'normalize', 'normalize_into', 'reformat', 'reformat_into', 'remap', 'remap_into', 'resize', 'resize_into', 'rotate', 'rotate_into', 'threshold', 'threshold_into', 'warp_affine', 'warp_affine_into', 'warp_perspective', 'warp_perspective_into'] +__all__: list[str] = ['BGR', 'BINARY', 'BINARY_INV', 'BOTH', 'BT2020', 'BT601', 'BT709', 'BndBox', 'BndBoxes', 'Box', 'CHW', 'COLOR_BGR2GRAY', 'COLOR_BGR2RGB', 'COLOR_BGR2YUV', 'COLOR_BGR2YUV_NV12', 'COLOR_BGR2YUV_NV21', 'COLOR_RGB2BGR', 'COLOR_RGB2GRAY', 'COLOR_RGB2YUV', 'COLOR_RGB2YUV_NV12', 'COLOR_RGB2YUV_NV21', 'COLOR_YUV2BGR', 'COLOR_YUV2BGR_NV12', 'COLOR_YUV2BGR_NV21', 'COLOR_YUV2RGB', 'COLOR_YUV2RGB_NV12', 'COLOR_YUV2RGB_NV21', 'CONSTANT', 'CPU', 'CUBIC', 'ColorRGBA', 'Exception', 'F32', 'F64', 'Format', 'GPU', 'Grayscale', 'HWC', 'Image', 'ImageBatchVarShape', 'LINEAR', 'N', 'NC', 'NCHW', 'NEAREST', 'NHWC', 'NW', 'NWC', 'NormalizeFlags', 'REFLECT', 'REFLECT101', 'REMAP_ABSOLUTE', 'REMAP_ABSOLUTE_NORMALIZED', 'REMAP_RELATIVE_NORMALIZED', 'REPLICATE', 'RGB', 'S16', 'S32', 'S8', 'Size2D', 'Stream', 'TOZERO', 'TOZERO_INV', 'TRUNC', 'Tensor', 'U16', 'U32', 'U8', 'WRAP', 'X', 'XYZW', 'Y', 'YUV', 'YVU', 'ZYXW', 'advcvtcolor', 'advcvtcolor_into', 'as_image', 'as_images', 'bilateral_filter', 'bilateral_filter_into', 'bndbox', 'bndbox_into', 'center_crop', 'center_crop_into', 'composite', 'composite_into', 'convert_to', 'convert_to_into', 'copymakeborder', 'copymakeborder_into', 'custom_crop', 'custom_crop_into', 'cvtcolor', 'cvtcolor_into', 'eAxis', 'eBorderType', 'eChannelType', 'eColorConversionCode', 'eColorSpec', 'eDataType', 'eDeviceType', 'eInterpolationType', 'eRemapType', 'eSwizzle', 'eTensorLayout', 'eThresholdType', 'flip', 'flip_into', 'from_dlpack', 'gamma_contrast', 'gamma_contrast_into', 'histogram', 'histogram_into', 'nms', 'nms_into', 'normalize', 'normalize_into', 'reformat', 'reformat_into', 'remap', 'remap_into', 'resize', 'resize_into', 'rotate', 'rotate_into', 'threshold', 'threshold_into', 'warp_affine', 'warp_affine_into', 'warp_perspective', 'warp_perspective_into'] class BndBox: borderColor: ColorRGBA box: Box @@ -92,6 +92,187 @@ class ColorRGBA: ... class Exception(Exception): pass +class _FormatConstants(type): + """ + Metaclass carrying the named Format constants. Declaring them here (rather + than on Format itself) makes them accessible on the class object only + (Format.RGB8) and not on Format *instances*, so chains like + Format.RGB8.RGB8 are correctly rejected by type checkers — mirroring the + runtime __getattribute__ guard in the C++ binding. + """ + BGR8: Format # value = rocpycv.Format.BGR8 + BGRA8: Format # value = rocpycv.Format.BGRA8 + F32: Format # value = rocpycv.Format.F32 + F64: Format # value = rocpycv.Format.F64 + NONE: Format # value = rocpycv.Format.NONE + RGB16: Format # value = rocpycv.Format.RGB16 + RGB32: Format # value = rocpycv.Format.RGB32 + RGB8: Format # value = rocpycv.Format.RGB8 + RGBA16: Format # value = rocpycv.Format.RGBA16 + RGBA32: Format # value = rocpycv.Format.RGBA32 + RGBA8: Format # value = rocpycv.Format.RGBA8 + RGBAf32: Format # value = rocpycv.Format.RGBAf32 + RGBAf64: Format # value = rocpycv.Format.RGBAf64 + RGBAs16: Format # value = rocpycv.Format.RGBAs16 + RGBAs32: Format # value = rocpycv.Format.RGBAs32 + RGBAs8: Format # value = rocpycv.Format.RGBAs8 + RGBf32: Format # value = rocpycv.Format.RGBf32 + RGBf64: Format # value = rocpycv.Format.RGBf64 + RGBs16: Format # value = rocpycv.Format.RGBs16 + RGBs32: Format # value = rocpycv.Format.RGBs32 + RGBs8: Format # value = rocpycv.Format.RGBs8 + S16: Format # value = rocpycv.Format.S16 + S32: Format # value = rocpycv.Format.S32 + S8: Format # value = rocpycv.Format.S8 + U16: Format # value = rocpycv.Format.U16 + U32: Format # value = rocpycv.Format.U32 + U8: Format # value = rocpycv.Format.U8 + _2F32: Format # value = rocpycv.Format._2F32 + +class Format(metaclass=_FormatConstants): + """ + Describes how image pixel data is laid out in memory. + """ + def __eq__(self, arg0: Format) -> bool: + ... + def __hash__(self) -> int: + ... + def __init__(self, dtype: eDataType, channels: typing.SupportsInt | typing.SupportsIndex, swizzle: eSwizzle = ...) -> None: + """ + Construct a format from a data type, channel count, and swizzle. + """ + def __ne__(self, arg0: Format) -> bool: + ... + def __repr__(self) -> str: + ... + @property + def channels(self) -> int: + """ + Read-only property that returns the number of color channels in the image. + """ + @property + def dtype(self) -> eDataType: + """ + Read-only property that returns the data type of each channel. + """ + @property + def planes(self) -> int: + """ + Read-only property that returns the number of planes in the image (always 1; rocCV is single-plane). + """ + @property + def swizzle(self) -> eSwizzle: + """ + Read-only property that returns the channel swizzle of the format. + """ +class Image: + """ + A single variable-sized image with device-resident data. + """ + @staticmethod + def zeros(size: tuple[typing.SupportsInt | typing.SupportsIndex, typing.SupportsInt | typing.SupportsIndex], format: Format, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Image: + """ + Create an image of the given size and format filled with zeros. + """ + def __dlpack__(self, stream: typing.Any = None) -> typing_extensions.CapsuleType: + """ + Creates a DLPack capsule from this image. + """ + def __dlpack_device__(self) -> tuple: + """ + Returns a tuple containing the DLPack device type and id for this image. + """ + def __init__(self, size: tuple[typing.SupportsInt | typing.SupportsIndex, typing.SupportsInt | typing.SupportsIndex], format: Format, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: + """ + Allocate a new image of the given size (width, height) and format. + """ + def copy_to(self, device: eDeviceType) -> Image: + """ + Returns a copy of the image with data copied to the given device. + """ + @property + def device(self) -> eDeviceType: + """ + Read-only property returning the device of the image. + """ + @property + def format(self) -> Format: + """ + Read-only property returning the format of the image. + """ + @property + def height(self) -> int: + """ + Read-only property returning the height of the image. + """ + @property + def size(self) -> tuple[int, int]: + """ + Read-only property returning the (width, height) of the image. + """ + @property + def width(self) -> int: + """ + Read-only property returning the width of the image. + """ +class ImageBatchVarShape: + """ + A batch of variable-sized images that may differ in size, format, and dtype. + """ + def __getitem__(self, index: typing.SupportsInt | typing.SupportsIndex) -> Image: + """ + Return the image at the given index. + """ + def __init__(self, capacity: typing.SupportsInt | typing.SupportsIndex, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: + """ + Create an empty batch with the given capacity (maximum number of images). + """ + def __iter__(self) -> collections.abc.Iterator[Image]: + """ + Return an iterator over the images in the batch. + """ + def __len__(self) -> int: + """ + Return the number of images currently in the batch. + """ + def clear(self) -> None: + """ + Remove all images from the batch; capacity is retained. + """ + def popback(self, count: typing.SupportsInt | typing.SupportsIndex = 1) -> None: + """ + Remove one or more images from the end of the batch. + """ + @typing.overload + def pushback(self, image: Image) -> None: + """ + Append a single image to the end of the batch. + """ + @typing.overload + def pushback(self, images: collections.abc.Sequence[Image]) -> None: + """ + Append multiple images to the end of the batch. + """ + @property + def capacity(self) -> int: + """ + Read-only property returning the maximum number of images the batch can hold. + """ + @property + def device(self) -> eDeviceType: + """ + Read-only property returning the device the batch resides on. + """ + @property + def maxsize(self) -> tuple[int, int]: + """ + Read-only property returning the (width, height) bounding box across all images. + """ + @property + def uniqueformat(self) -> Format: + """ + Read-only property returning the common format, or Format.NONE if heterogeneous/empty. + """ class NormalizeFlags: """ Members: @@ -166,7 +347,7 @@ class Tensor: """ Returns a tuple containing the DLPack device and device id for the tensor. """ - def __init__(self, shape: collections.abc.Sequence[typing.SupportsInt | typing.SupportsIndex], layout: eTensorLayout, dtype: eDataType, device: eDeviceType = ...) -> None: + def __init__(self, shape: collections.abc.Sequence[typing.SupportsInt | typing.SupportsIndex], layout: eTensorLayout, dtype: eDataType, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Constructs a tensor object. """ @@ -627,6 +808,43 @@ class eRemapType: @property def value(self) -> int: ... +class eSwizzle: + """ + Members: + + XYZW + + ZYXW + """ + XYZW: typing.ClassVar[eSwizzle] # value = + ZYXW: typing.ClassVar[eSwizzle] # value = + __members__: typing.ClassVar[dict[str, eSwizzle]] # value = {'XYZW': , 'ZYXW': } + def __eq__(self, other: typing.Any) -> bool: + ... + def __getstate__(self) -> int: + ... + def __hash__(self) -> int: + ... + def __index__(self) -> int: + ... + def __init__(self, value: typing.SupportsInt | typing.SupportsIndex) -> None: + ... + def __int__(self) -> int: + ... + def __ne__(self, other: typing.Any) -> bool: + ... + def __repr__(self) -> str: + ... + def __setstate__(self, state: typing.SupportsInt | typing.SupportsIndex) -> None: + ... + def __str__(self) -> str: + ... + @property + def name(self) -> str: + ... + @property + def value(self) -> int: + ... class eTensorLayout: """ Members: @@ -728,7 +946,7 @@ class eThresholdType: @property def value(self) -> int: ... -def advcvtcolor(src: Tensor, conversion_code: eColorConversionCode, color_spec: eColorSpec, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def advcvtcolor(src: Tensor, conversion_code: eColorConversionCode, color_spec: eColorSpec, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Advanced Color Convert operation on the given HIP stream. @@ -745,7 +963,7 @@ def advcvtcolor(src: Tensor, conversion_code: eColorConversionCode, color_spec: Returns: rocpycv.Tensor: The output tensor. """ -def advcvtcolor_into(dst: Tensor, src: Tensor, conversion_code: eColorConversionCode, color_spec: eColorSpec, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def advcvtcolor_into(dst: Tensor, src: Tensor, conversion_code: eColorConversionCode, color_spec: eColorSpec, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Advanced Color Convert operation on the given HIP stream. @@ -763,7 +981,15 @@ def advcvtcolor_into(dst: Tensor, src: Tensor, conversion_code: eColorConversion Returns: None """ -def bilateral_filter(src: Tensor, diameter: typing.SupportsInt | typing.SupportsIndex, sigmaColor: typing.SupportsFloat | typing.SupportsIndex, sigmaSpace: typing.SupportsFloat | typing.SupportsIndex, borderMode: eBorderType, borderValue: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def as_image(buffer: typing.Any, format: Format = ...) -> Image: + """ + Wraps a DLPack-supported buffer as a rocpycv Image without copying. + """ +def as_images(buffers: collections.abc.Sequence[typing.Any], format: Format = ...) -> ImageBatchVarShape: + """ + Wraps a list of DLPack-supported buffers as a rocpycv ImageBatchVarShape without copying. + """ +def bilateral_filter(src: Tensor, diameter: typing.SupportsInt | typing.SupportsIndex, sigmaColor: typing.SupportsFloat | typing.SupportsIndex, sigmaSpace: typing.SupportsFloat | typing.SupportsIndex, borderMode: eBorderType, borderValue: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Bilateral Filter operation on the given HIP stream. @@ -783,7 +1009,7 @@ def bilateral_filter(src: Tensor, diameter: typing.SupportsInt | typing.Supports Returns: rocpycv.Tensor: The output tensor. """ -def bilateral_filter_into(dst: Tensor, src: Tensor, diameter: typing.SupportsInt | typing.SupportsIndex, sigmaColor: typing.SupportsFloat | typing.SupportsIndex, sigmaSpace: typing.SupportsFloat | typing.SupportsIndex, borderMode: eBorderType, borderValue: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def bilateral_filter_into(dst: Tensor, src: Tensor, diameter: typing.SupportsInt | typing.SupportsIndex, sigmaColor: typing.SupportsFloat | typing.SupportsIndex, sigmaSpace: typing.SupportsFloat | typing.SupportsIndex, borderMode: eBorderType, borderValue: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Bilateral Filter operation on the given HIP stream. @@ -804,7 +1030,7 @@ def bilateral_filter_into(dst: Tensor, src: Tensor, diameter: typing.SupportsInt Returns: None """ -def bndbox(src: Tensor, bnd_boxes: BndBoxes, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def bndbox(src: Tensor, bnd_boxes: BndBoxes, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the BndBox operation on the given HIP stream. @@ -820,7 +1046,7 @@ def bndbox(src: Tensor, bnd_boxes: BndBoxes, stream: rocpycv.Stream | None = Non Returns: rocpycv.Tensor: The output tensor. """ -def bndbox_into(dst: Tensor, src: Tensor, bnd_boxes: BndBoxes, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def bndbox_into(dst: Tensor, src: Tensor, bnd_boxes: BndBoxes, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the BndBox operation on the given HIP stream. @@ -837,7 +1063,7 @@ def bndbox_into(dst: Tensor, src: Tensor, bnd_boxes: BndBoxes, stream: rocpycv.S Returns: None """ -def center_crop(src: Tensor, crop_size: tuple, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def center_crop(src: Tensor, crop_size: tuple, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Center Crop operation on the given HIP stream. @@ -854,7 +1080,7 @@ def center_crop(src: Tensor, crop_size: tuple, stream: rocpycv.Stream | None = N Returns: rocpycv.Tensor: The output tensor. """ -def center_crop_into(dst: Tensor, src: Tensor, crop_size: tuple, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def center_crop_into(dst: Tensor, src: Tensor, crop_size: tuple, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Center Crop operation on the given HIP stream. @@ -871,7 +1097,7 @@ def center_crop_into(dst: Tensor, src: Tensor, crop_size: tuple, stream: rocpycv Returns: None """ -def composite(foreground: Tensor, background: Tensor, fgmask: Tensor, outchannels: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def composite(foreground: Tensor, background: Tensor, fgmask: Tensor, outchannels: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Composite operation on the given HIP stream. @@ -889,7 +1115,7 @@ def composite(foreground: Tensor, background: Tensor, fgmask: Tensor, outchannel Returns: rocpycv.Tensor: The output tensor with number of channels. """ -def composite_into(dst: Tensor, foreground: Tensor, background: Tensor, fgmask: Tensor, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def composite_into(dst: Tensor, foreground: Tensor, background: Tensor, fgmask: Tensor, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Composite operation on the given HIP stream. @@ -907,7 +1133,7 @@ def composite_into(dst: Tensor, foreground: Tensor, background: Tensor, fgmask: Returns: None """ -def convert_to(src: Tensor, dtype: eDataType, alpha: typing.SupportsFloat | typing.SupportsIndex = 1.0, beta: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def convert_to(src: Tensor, dtype: eDataType, alpha: typing.SupportsFloat | typing.SupportsIndex = 1.0, beta: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Convert To operation on the given HIP stream. @@ -925,7 +1151,7 @@ def convert_to(src: Tensor, dtype: eDataType, alpha: typing.SupportsFloat | typi Returns: rocpycv.Tensor: The output tensor. """ -def convert_to_into(dst: Tensor, src: Tensor, alpha: typing.SupportsFloat | typing.SupportsIndex = 1.0, beta: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def convert_to_into(dst: Tensor, src: Tensor, alpha: typing.SupportsFloat | typing.SupportsIndex = 1.0, beta: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Convert To operation on the given HIP stream. @@ -943,7 +1169,7 @@ def convert_to_into(dst: Tensor, src: Tensor, alpha: typing.SupportsFloat | typi Returns: None """ -def copymakeborder(src: Tensor, border_mode: eBorderType = ..., border_value: list = [0.0, 0.0, 0.0, 0.0], top: typing.SupportsInt | typing.SupportsIndex, bottom: typing.SupportsInt | typing.SupportsIndex, left: typing.SupportsInt | typing.SupportsIndex, right: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def copymakeborder(src: Tensor, border_mode: eBorderType = eBorderType.eBorderType.CONSTANT, border_value: list = [0.0, 0.0, 0.0, 0.0], top: typing.SupportsInt | typing.SupportsIndex, bottom: typing.SupportsInt | typing.SupportsIndex, left: typing.SupportsInt | typing.SupportsIndex, right: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the CopyMakeBorder operation on the given HIP stream. @@ -964,7 +1190,7 @@ def copymakeborder(src: Tensor, border_mode: eBorderType = ..., border_value: li Returns: rocpycv.Tensor: The output tensor. """ -def copymakeborder_into(dst: Tensor, src: Tensor, border_mode: eBorderType = ..., border_value: list = [0.0, 0.0, 0.0, 0.0], top: typing.SupportsInt | typing.SupportsIndex, left: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def copymakeborder_into(dst: Tensor, src: Tensor, border_mode: eBorderType = eBorderType.eBorderType.CONSTANT, border_value: list = [0.0, 0.0, 0.0, 0.0], top: typing.SupportsInt | typing.SupportsIndex, left: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the CopyMakeBorder operation on the given HIP stream. @@ -984,7 +1210,7 @@ def copymakeborder_into(dst: Tensor, src: Tensor, border_mode: eBorderType = ... Returns: None """ -def custom_crop(src: Tensor, crop_rect: Box, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def custom_crop(src: Tensor, crop_rect: Box, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Custom Crop operation on the given HIP stream. @@ -1001,7 +1227,7 @@ def custom_crop(src: Tensor, crop_rect: Box, stream: rocpycv.Stream | None = Non Returns: None """ -def custom_crop_into(dst: Tensor, src: Tensor, crop_rect: Box, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def custom_crop_into(dst: Tensor, src: Tensor, crop_rect: Box, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Custom Crop operation on the given HIP stream. @@ -1017,7 +1243,7 @@ def custom_crop_into(dst: Tensor, src: Tensor, crop_rect: Box, stream: rocpycv.S Returns: rocpycv.Tensor: The output tensor. """ -def cvtcolor(src: Tensor, conversion_code: eColorConversionCode, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def cvtcolor(src: Tensor, conversion_code: eColorConversionCode, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Color Convert operation on the given HIP stream. @@ -1033,7 +1259,7 @@ def cvtcolor(src: Tensor, conversion_code: eColorConversionCode, stream: rocpycv Returns: rocpycv.Tensor: The output tensor. """ -def cvtcolor_into(dst: Tensor, src: Tensor, conversion_code: eColorConversionCode, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def cvtcolor_into(dst: Tensor, src: Tensor, conversion_code: eColorConversionCode, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Color Convert operation on the given HIP stream. @@ -1050,7 +1276,7 @@ def cvtcolor_into(dst: Tensor, src: Tensor, conversion_code: eColorConversionCod Returns: None """ -def flip(src: Tensor, flip_code: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def flip(src: Tensor, flip_code: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Flip operation on the given HIP stream. @@ -1066,7 +1292,7 @@ def flip(src: Tensor, flip_code: typing.SupportsInt | typing.SupportsIndex, stre Returns: rocpycv.Tensor: The output tensor. """ -def flip_into(dst: Tensor, src: Tensor, flip_code: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def flip_into(dst: Tensor, src: Tensor, flip_code: typing.SupportsInt | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Flip operation on the given HIP stream. @@ -1087,7 +1313,7 @@ def from_dlpack(buffer: typing.Any, layout: eTensorLayout) -> Tensor: """ Wraps a DLPack supported tensor in a rocpycv tensor. """ -def gamma_contrast(src: Tensor, gamma: typing.SupportsFloat | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def gamma_contrast(src: Tensor, gamma: typing.SupportsFloat | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Gamma Contrast operation on the given HIP stream. @@ -1103,7 +1329,7 @@ def gamma_contrast(src: Tensor, gamma: typing.SupportsFloat | typing.SupportsInd Returns: rocpycv.Tensor: The output tensor. """ -def gamma_contrast_into(dst: Tensor, src: Tensor, gamma: typing.SupportsFloat | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def gamma_contrast_into(dst: Tensor, src: Tensor, gamma: typing.SupportsFloat | typing.SupportsIndex, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Gamma Contrast operation on the given HIP stream. @@ -1120,7 +1346,7 @@ def gamma_contrast_into(dst: Tensor, src: Tensor, gamma: typing.SupportsFloat | Returns: None """ -def histogram(src: Tensor, mask: rocpycv.Tensor | None, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def histogram(src: Tensor, mask: rocpycv.Tensor | None, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Histogram operation on the given HIP stream. @@ -1136,7 +1362,7 @@ def histogram(src: Tensor, mask: rocpycv.Tensor | None, stream: rocpycv.Stream | Returns: rocpycv.Tensor: Output tensor with width of 256 and a height equal to the batch size of input (1 if HWC input). """ -def histogram_into(dst: Tensor, src: Tensor, mask: rocpycv.Tensor | None, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def histogram_into(dst: Tensor, src: Tensor, mask: rocpycv.Tensor | None, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Histogram operation on the given HIP stream. @@ -1153,7 +1379,7 @@ def histogram_into(dst: Tensor, src: Tensor, mask: rocpycv.Tensor | None, stream Returns: None """ -def nms(src: Tensor, scores: Tensor, score_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.1920928955078125e-07, iou_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.0, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def nms(src: Tensor, scores: Tensor, score_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.1920928955078125e-07, iou_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.0, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Non-maximum Suppression operation on the given HIP stream. @@ -1171,7 +1397,7 @@ def nms(src: Tensor, scores: Tensor, score_threshold: typing.SupportsFloat | typ Returns: rocpycv.Tensor: The output tensor of shape [i, j], containing 1 (kept) or 0 (suppressed) for each bounding box (j) per batch (i). Results will be written to this tensor. """ -def nms_into(dst: Tensor, src: Tensor, scores: Tensor, score_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.1920928955078125e-07, iou_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.0, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def nms_into(dst: Tensor, src: Tensor, scores: Tensor, score_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.1920928955078125e-07, iou_threshold: typing.SupportsFloat | typing.SupportsIndex = 1.0, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Non-maximum Suppression operation on the given HIP stream. @@ -1190,7 +1416,7 @@ def nms_into(dst: Tensor, src: Tensor, scores: Tensor, score_threshold: typing.S Returns: None """ -def normalize(src: Tensor, base: Tensor, scale: Tensor, flags: typing.SupportsInt | typing.SupportsIndex | None = None, globalscale: typing.SupportsFloat | typing.SupportsIndex = 1.0, globalshift: typing.SupportsFloat | typing.SupportsIndex = 0.0, epsilon: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def normalize(src: Tensor, base: Tensor, scale: Tensor, flags: typing.SupportsInt | typing.SupportsIndex | None = None, globalscale: typing.SupportsFloat | typing.SupportsIndex = 1.0, globalshift: typing.SupportsFloat | typing.SupportsIndex = 0.0, epsilon: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Normalize operation on the given HIP stream. @@ -1211,7 +1437,7 @@ def normalize(src: Tensor, base: Tensor, scale: Tensor, flags: typing.SupportsIn Returns: rocpycv.Tensor: The output tensor. """ -def normalize_into(dst: Tensor, src: Tensor, base: Tensor, scale: Tensor, flags: typing.SupportsInt | typing.SupportsIndex | None = None, globalscale: typing.SupportsFloat | typing.SupportsIndex = 1.0, globalshift: typing.SupportsFloat | typing.SupportsIndex = 0.0, epsilon: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def normalize_into(dst: Tensor, src: Tensor, base: Tensor, scale: Tensor, flags: typing.SupportsInt | typing.SupportsIndex | None = None, globalscale: typing.SupportsFloat | typing.SupportsIndex = 1.0, globalshift: typing.SupportsFloat | typing.SupportsIndex = 0.0, epsilon: typing.SupportsFloat | typing.SupportsIndex = 0.0, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Normalize operation on the given HIP stream. @@ -1233,7 +1459,7 @@ def normalize_into(dst: Tensor, src: Tensor, base: Tensor, scale: Tensor, flags: Returns: None """ -def reformat(input: Tensor, out_layout: eTensorLayout, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def reformat(input: Tensor, out_layout: eTensorLayout, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Reformat operation and returns the result as a new tensor. @@ -1249,7 +1475,7 @@ def reformat(input: Tensor, out_layout: eTensorLayout, stream: rocpycv.Stream | Returns: rocpycv.Tensor: The reformatted tensor. """ -def reformat_into(output: Tensor, input: Tensor, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def reformat_into(output: Tensor, input: Tensor, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Reformat operation on the given HIP stream. @@ -1265,7 +1491,7 @@ def reformat_into(output: Tensor, input: Tensor, stream: rocpycv.Stream | None = Returns: None """ -def remap(src: Tensor, map: Tensor, in_interpolation: eInterpolationType, map_interpolation: eInterpolationType, map_value_type: eRemapType, align_corners: bool, border_type: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def remap(src: Tensor, map: Tensor, in_interpolation: eInterpolationType, map_interpolation: eInterpolationType, map_value_type: eRemapType, align_corners: bool, border_type: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Remap operation on the given HIP stream. @@ -1287,7 +1513,7 @@ def remap(src: Tensor, map: Tensor, in_interpolation: eInterpolationType, map_in Returns: rocpycv.Tensor: The output tensor. """ -def remap_into(dst: Tensor, src: Tensor, map: Tensor, in_interpolation: eInterpolationType, map_interpolation: eInterpolationType, map_value_type: eRemapType, align_corners: bool, border_type: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def remap_into(dst: Tensor, src: Tensor, map: Tensor, in_interpolation: eInterpolationType, map_interpolation: eInterpolationType, map_value_type: eRemapType, align_corners: bool, border_type: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Remap operation on the given HIP stream. @@ -1310,7 +1536,7 @@ def remap_into(dst: Tensor, src: Tensor, map: Tensor, in_interpolation: eInterpo Returns: None """ -def resize(src: Tensor, shape: tuple, interp: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def resize(src: Tensor, shape: tuple, interp: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Resize operation on the given HIP stream. @@ -1327,7 +1553,7 @@ def resize(src: Tensor, shape: tuple, interp: eInterpolationType, stream: rocpyc Returns: rocpycv.Tensor: The output tensor. """ -def resize_into(dst: Tensor, src: Tensor, interp: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def resize_into(dst: Tensor, src: Tensor, interp: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Resize operation on the given HIP stream. @@ -1344,7 +1570,7 @@ def resize_into(dst: Tensor, src: Tensor, interp: eInterpolationType, stream: ro Returns: None """ -def rotate(src: Tensor, angle_deg: typing.SupportsFloat | typing.SupportsIndex, shift: tuple, interpolation: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def rotate(src: Tensor, angle_deg: typing.SupportsFloat | typing.SupportsIndex, shift: tuple, interpolation: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Rotate operation on the given HIP stream. @@ -1362,7 +1588,7 @@ def rotate(src: Tensor, angle_deg: typing.SupportsFloat | typing.SupportsIndex, Returns: rocpycv.Tensor: The output tensor. """ -def rotate_into(dst: Tensor, src: Tensor, angle_deg: typing.SupportsFloat | typing.SupportsIndex, shift: tuple, interpolation: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def rotate_into(dst: Tensor, src: Tensor, angle_deg: typing.SupportsFloat | typing.SupportsIndex, shift: tuple, interpolation: eInterpolationType, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Rotate operation on the given HIP stream. @@ -1381,7 +1607,7 @@ def rotate_into(dst: Tensor, src: Tensor, angle_deg: typing.SupportsFloat | typi Returns: None """ -def threshold(src: Tensor, thresh: Tensor, maxVal: Tensor, maxBatchSize: typing.SupportsInt | typing.SupportsIndex, threshType: eThresholdType, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def threshold(src: Tensor, thresh: Tensor, maxVal: Tensor, maxBatchSize: typing.SupportsInt | typing.SupportsIndex, threshType: eThresholdType, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Thresholding operation on the given HIP stream. @@ -1397,7 +1623,7 @@ def threshold(src: Tensor, thresh: Tensor, maxVal: Tensor, maxBatchSize: typing. stream (rocpycv.Stream, optional): HIP stream to run this operation on. device (rocpycv.Device, optional): The device to run this operation on. Defaults to GPU. """ -def threshold_into(dst: Tensor, src: Tensor, thresh: Tensor, maxVal: Tensor, maxBatchSize: typing.SupportsInt | typing.SupportsIndex, threshType: eThresholdType, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def threshold_into(dst: Tensor, src: Tensor, thresh: Tensor, maxVal: Tensor, maxBatchSize: typing.SupportsInt | typing.SupportsIndex, threshType: eThresholdType, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Thresholding operation on the given HIP stream. @@ -1414,7 +1640,7 @@ def threshold_into(dst: Tensor, src: Tensor, thresh: Tensor, maxVal: Tensor, max stream (rocpycv.Stream, optional): HIP stream to run this operation on. device (rocpycv.Device, optional): The device to run this operation on. Defaults to GPU. """ -def warp_affine(src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def warp_affine(src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Warp Affine operation on the given HIP stream. @@ -1434,7 +1660,7 @@ def warp_affine(src: Tensor, xform: list, inverted: bool, interp: eInterpolation Returns: rocpycv.Tensor: The output tensor. """ -def warp_affine_into(dst: Tensor, src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def warp_affine_into(dst: Tensor, src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Warp Affine operation on the given HIP stream. @@ -1455,7 +1681,7 @@ def warp_affine_into(dst: Tensor, src: Tensor, xform: list, inverted: bool, inte Returns: None """ -def warp_perspective(src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> Tensor: +def warp_perspective(src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> Tensor: """ Executes the Warp Perspective operation on the given HIP stream. @@ -1475,7 +1701,7 @@ def warp_perspective(src: Tensor, xform: list, inverted: bool, interp: eInterpol Returns: rocpycv.Tensor: The output tensor. """ -def warp_perspective_into(dst: Tensor, src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = ...) -> None: +def warp_perspective_into(dst: Tensor, src: Tensor, xform: list, inverted: bool, interp: eInterpolationType, border_mode: eBorderType, border_value: list, stream: rocpycv.Stream | None = None, device: eDeviceType = eDeviceType.eDeviceType.GPU) -> None: """ Executes the Warp Perspective operation on the given HIP stream. @@ -1554,6 +1780,8 @@ U32: eDataType # value = U8: eDataType # value = WRAP: eBorderType # value = X: eAxis # value = +XYZW: eSwizzle # value = Y: eAxis # value = YUV: eChannelType # value = YVU: eChannelType # value = +ZYXW: eSwizzle # value = diff --git a/samples/CMakeLists.txt b/samples/CMakeLists.txt index a1a06a30..0a2b9d83 100644 --- a/samples/CMakeLists.txt +++ b/samples/CMakeLists.txt @@ -128,6 +128,7 @@ if (BUILD_SAMPLES) define_sample_executable(warp_perspective) define_sample_executable(composite) define_sample_executable(gamma_contrast) + define_sample_executable(resize_var_shape) add_subdirectory(cropandresize/cpp) else() message("-- ${Yellow}roccv-samples -- missing required dependencies, will not be built${ColorReset}") diff --git a/samples/README.md b/samples/README.md index fcbb98e6..0aa00292 100644 --- a/samples/README.md +++ b/samples/README.md @@ -13,7 +13,8 @@ apt install python3-opencv # For Python samples ## Operator Samples 1. Individual process in C++: bnd_box.cpp, center_crop.cpp, composite.cpp, copy_make_border.cpp, custom_crop.cpp, gamma_contrast.cpp, normalize.cpp, warp_perspective.cpp. 2. cropandresize - Crops and resizes the input image. This sample is designed to demonstrate a simple pipeline for multiple operators. -3. pipeline/multi_op_1.py: A pipeline of color conversion, cropping, bilateral filtering, bounding box drawing, rotation and resizing. +3. resize_var_shape.cpp - Loads a directory of variably-sized images into an ImageBatchVarShape and resizes the whole batch into a single uniform, constant-sized output tensor (GPU only). Output size is configurable via `-W`/`-H`. +4. pipeline/multi_op_1.py: A pipeline of color conversion, cropping, bilateral filtering, bounding box drawing, rotation and resizing. ## Building and running the samples Build rocCV as described in the main README with the `-D SAMPLES=ON` flag set. diff --git a/samples/resize_var_shape.cpp b/samples/resize_var_shape.cpp new file mode 100644 index 00000000..521e10be --- /dev/null +++ b/samples/resize_var_shape.cpp @@ -0,0 +1,196 @@ +/* + * Copyright (c) 2025 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 +#include +#include +#include +#include + +#include "common/utils.hpp" + +/** + * @brief Variable-shape Resize sample app. + * + * Loads every image in a directory (each potentially a different size) into a single + * roccv::ImageBatchVarShape, then resizes the whole batch into one uniform, constant-sized + * output tensor using the variable-shape Resize overload. This is a GPU-only path. + * + * Image Directory -> ImageBatchVarShape -> Resize (varshape) -> WriteImage + */ + +void ShowUsage() { + std::cout << "usage: ./resize_var_shape -i [-W ] [-H ]\n" + << " -i Path to a directory of images (or a single image file). Required.\n" + << " -W Output width in pixels. Optional; default: 224.\n" + << " -H Output height in pixels. Optional; default: 224.\n" + << " -h Show this help message.\n" + << "Resized images are written to ./output/image_.bmp.\n"; +} + +int ParseArgs(int argc, char *argv[], std::string &imageDir, int &outWidth, int &outHeight) { + static struct option long_options[] = {{"help", no_argument, 0, 'h'}, + {"imageDir", required_argument, 0, 'i'}, + {"width", required_argument, 0, 'W'}, + {"height", required_argument, 0, 'H'}, + {0, 0, 0, 0}}; + + int long_index = 0; + int opt = 0; + while ((opt = getopt_long(argc, argv, "hi:W:H:", long_options, &long_index)) != -1) { + switch (opt) { + case 'h': + ShowUsage(); + return -1; + case 'i': + imageDir = optarg; + break; + case 'W': + outWidth = std::stoi(optarg); + break; + case 'H': + outHeight = std::stoi(optarg); + break; + default: + ShowUsage(); + return -1; + } + } + + if (imageDir.empty()) { + ShowUsage(); + std::cerr << "Error: an input image directory (-i) is required.\n"; + return -1; + } + if (outWidth <= 0 || outHeight <= 0) { + ShowUsage(); + std::cerr << "Error: output width/height must be positive.\n"; + return -1; + } + return 0; +} + +/** + * @brief Collects supported image files from a directory (non-recursive) or accepts a single image file. + */ +std::vector CollectImageFiles(const std::string &path) { + const std::vector supportedExtensions = {".bmp", ".jpg", ".jpeg", ".png"}; + + std::vector imageFiles; + if (std::filesystem::is_directory(path)) { + for (const auto &entry : std::filesystem::directory_iterator(path)) { + if (!std::filesystem::is_directory(entry.path()) && ContainsExtension(entry.path(), supportedExtensions)) { + imageFiles.push_back(entry.path().string()); + } + } + // Sort for a deterministic batch order across runs. + std::sort(imageFiles.begin(), imageFiles.end()); + } else { + if (!ContainsExtension(path, supportedExtensions)) { + throw std::runtime_error("Cannot decode " + path + ". File type not supported."); + } + imageFiles.push_back(path); + } + + if (imageFiles.empty()) { + throw std::runtime_error("No supported images found in " + path); + } + return imageFiles; +} + +int main(int argc, char *argv[]) { + std::string imageDir; + int outWidth = 224; + int outHeight = 224; + + int retval = ParseArgs(argc, argv, imageDir, outWidth, outHeight); + if (retval != 0) { + return retval; + } + + try { + // tag: Create the HIP stream + hipStream_t stream; + CHECK_HIP_ERROR(hipStreamCreate(&stream)); + + // tag: Gather input images + std::vector imageFiles = CollectImageFiles(imageDir); + const int numImages = static_cast(imageFiles.size()); + std::cout << "Found " << numImages << " image(s) in " << imageDir << ".\n"; + + // tag: Build a variable-shape image batch + // Each image keeps its own (potentially unique) dimensions. The batch holds GPU-resident images; the + // input is loaded host-side via OpenCV and copied straight into each image's device plane. + roccv::ImageBatchVarShape batch(numImages); + + for (const std::string &file : imageFiles) { + // OpenCV loads as 8-bit BGR; we treat the interleaved 3-channel buffer as FMT_RGB8 without a color + // conversion (output is written back through OpenCV, so the channel order round-trips consistently). + cv::Mat inputMat = cv::imread(file, cv::IMREAD_COLOR); + if (inputMat.empty()) { + throw std::runtime_error("Unable to load image " + file); + } + + roccv::Size2D size{inputMat.cols, inputMat.rows}; + roccv::Image image(size, roccv::FMT_RGB8, eDeviceType::GPU); + + // Copy host pixels into the image's GPU plane, honoring both the source (cv::Mat) and destination + // (rocCV plane) row strides. + auto imageData = image.exportData(); + const roccv::ImagePlaneStrided &plane = imageData.plane(0); + const size_t rowBytes = static_cast(inputMat.cols) * inputMat.channels() * sizeof(uint8_t); + CHECK_HIP_ERROR(hipMemcpy2D(plane.basePtr, plane.rowStride, inputMat.data, inputMat.step, rowBytes, + inputMat.rows, hipMemcpyHostToDevice)); + + batch.pushBack(image); + std::cout << " loaded " << file << " (" << size.w << "x" << size.h << ")\n"; + } + + // tag: Allocate the uniform, constant-sized output tensor (NHWC, GPU) + roccv::Tensor outputTensor(numImages, {outWidth, outHeight}, roccv::FMT_RGB8); + + // tag: Run the variable-shape Resize. Every image in the batch is resized into the same output size. + std::cout << "Resizing batch to " << outWidth << "x" << outHeight << " (LINEAR interpolation)...\n"; + roccv::Resize resizeOp; + resizeOp(stream, batch, outputTensor, INTERP_TYPE_LINEAR, eDeviceType::GPU); + + // tag: Copy results back to the host and write each resized image to ./output/image_.bmp + WriteImages(stream, outputTensor, "./output"); + std::cout << "Wrote " << numImages << " resized image(s) to ./output/image_.bmp\n"; + + // tag: Clean up + CHECK_HIP_ERROR(hipStreamDestroy(stream)); + } catch (const std::exception &e) { + std::cerr << "Error: " << e.what() << std::endl; + return 1; + } + + return 0; +} diff --git a/src/core/tensor_storage.cpp b/src/core/tensor_storage.cpp index 800f074e..6a0f7859 100644 --- a/src/core/tensor_storage.cpp +++ b/src/core/tensor_storage.cpp @@ -22,7 +22,6 @@ #include "core/tensor_storage.hpp" #include "core/detail/context.hpp" -#include "core/hip_assert.h" namespace roccv { TensorStorage::TensorStorage(void* data, eDeviceType device, eOwnership ownership) diff --git a/src/op_adv_cvt_color.cpp b/src/op_adv_cvt_color.cpp index 7f035507..c2c69175 100644 --- a/src/op_adv_cvt_color.cpp +++ b/src/op_adv_cvt_color.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/tensor.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/common/adv_cvt_color_coefficients.hpp" #include "kernels/device/adv_cvt_color_device.hpp" #include "kernels/host/adv_cvt_color_host.hpp" @@ -189,22 +189,22 @@ void AdvCvtColor::operator()(hipStream_t stream, const Tensor &input, Tensor &ou switch (conversionCode) { case COLOR_BGR2YUV: Kernels::Device::rgb_or_bgr_to_yuv_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta); break; case COLOR_RGB2YUV: Kernels::Device::rgb_or_bgr_to_yuv_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta); break; case COLOR_YUV2BGR: Kernels::Device::yuv_to_rgb_or_bgr_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta); break; case COLOR_YUV2RGB: Kernels::Device::yuv_to_rgb_or_bgr_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta); break; default: throw Exception("Unsupported conversion code.", eStatusType::INVALID_COMBINATION); @@ -215,26 +215,26 @@ void AdvCvtColor::operator()(hipStream_t stream, const Tensor &input, Tensor &ou if (outChannels == 3) { if (bgr) { - Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar3> - <<>>(ImageWrapper(input), ImageWrapper(output), + Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar3> + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { - Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar3> - <<>>(ImageWrapper(input), ImageWrapper(output), + Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar3> + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } else { if (bgr) { - Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar4> - <<>>(ImageWrapper(input), ImageWrapper(output), + Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar4> + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { - Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar4> - <<>>(ImageWrapper(input), ImageWrapper(output), + Kernels::Device::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar4> + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } @@ -245,21 +245,21 @@ void AdvCvtColor::operator()(hipStream_t stream, const Tensor &input, Tensor &ou if (inChannels == 3) { if (bgr) { Kernels::Device::rgb_or_bgr_to_nv12_or_nv21_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { Kernels::Device::rgb_or_bgr_to_nv12_or_nv21_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } else { if (bgr) { Kernels::Device::rgb_or_bgr_to_nv12_or_nv21_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { Kernels::Device::rgb_or_bgr_to_nv12_or_nv21_adv - <<>>(ImageWrapper(input), ImageWrapper(output), + <<>>(TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } @@ -268,23 +268,23 @@ void AdvCvtColor::operator()(hipStream_t stream, const Tensor &input, Tensor &ou if (IsInterleaved444(conversionCode)) { switch (conversionCode) { case COLOR_BGR2YUV: - Kernels::Host::rgb_or_bgr_to_yuv_adv(ImageWrapper(input), - ImageWrapper(output), coeff, + Kernels::Host::rgb_or_bgr_to_yuv_adv(TensorWrapper(input), + TensorWrapper(output), coeff, kDelta); break; case COLOR_RGB2YUV: - Kernels::Host::rgb_or_bgr_to_yuv_adv(ImageWrapper(input), - ImageWrapper(output), coeff, + Kernels::Host::rgb_or_bgr_to_yuv_adv(TensorWrapper(input), + TensorWrapper(output), coeff, kDelta); break; case COLOR_YUV2BGR: - Kernels::Host::yuv_to_rgb_or_bgr_adv(ImageWrapper(input), - ImageWrapper(output), coeff, + Kernels::Host::yuv_to_rgb_or_bgr_adv(TensorWrapper(input), + TensorWrapper(output), coeff, kDelta); break; case COLOR_YUV2RGB: - Kernels::Host::yuv_to_rgb_or_bgr_adv(ImageWrapper(input), - ImageWrapper(output), coeff, + Kernels::Host::yuv_to_rgb_or_bgr_adv(TensorWrapper(input), + TensorWrapper(output), coeff, kDelta); break; default: throw Exception("Unsupported conversion code.", eStatusType::INVALID_COMBINATION); @@ -292,41 +292,41 @@ void AdvCvtColor::operator()(hipStream_t stream, const Tensor &input, Tensor &ou } else if (IsSemiPlanarToInterleaved(conversionCode)) { if (outChannels == 3) { if (bgr) { - Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar3>( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar3>( + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { - Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar3>( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar3>( + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } else { if (bgr) { - Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar4>( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar4>( + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { - Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, - ImageWrapper, uchar4>( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + Kernels::Host::nv12_or_nv21_to_rgb_or_bgr_adv, + TensorWrapper, uchar4>( + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } } else { if (inChannels == 3) { if (bgr) { Kernels::Host::rgb_or_bgr_to_nv12_or_nv21_adv( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { Kernels::Host::rgb_or_bgr_to_nv12_or_nv21_adv( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } else { if (bgr) { Kernels::Host::rgb_or_bgr_to_nv12_or_nv21_adv( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } else { Kernels::Host::rgb_or_bgr_to_nv12_or_nv21_adv( - ImageWrapper(input), ImageWrapper(output), coeff, kDelta, uidx); + TensorWrapper(input), TensorWrapper(output), coeff, kDelta, uidx); } } } diff --git a/src/op_bilateral_filter.cpp b/src/op_bilateral_filter.cpp index dffba8ae..5c079b31 100644 --- a/src/op_bilateral_filter.cpp +++ b/src/op_bilateral_filter.cpp @@ -31,7 +31,7 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/detail/casting.hpp" #include "core/wrappers/border_wrapper.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/bilateral_filter_device.hpp" #include "kernels/host/bilateral_filter_host.hpp" @@ -43,8 +43,8 @@ BilateralFilter::~BilateralFilter() {} template void dispatch_bilateral_filter_border_mode(hipStream_t stream, const Tensor &input, const Tensor &output, int diameter, float sigmaColor, float sigmaSpace, T borderValue, eDeviceType device) { - BorderWrapper inputWrapper(input, borderValue); - ImageWrapper outputWrapper(output); + auto inputWrapper = MakeBorderWrapper(TensorWrapper(input), borderValue); + TensorWrapper outputWrapper(output); if (outputWrapper.channels() > 4 || outputWrapper.channels() < 1) { throw Exception("Invalid channel size: cannot be greater than 4 or less than 1.", eStatusType::OUT_OF_BOUNDS); @@ -61,8 +61,7 @@ void dispatch_bilateral_filter_border_mode(hipStream_t stream, const Tensor &inp sigmaSpace = 1.0f; } - const int radius = - (diameter <= 0) ? static_cast(std::roundf(sigmaSpace * 1.5f)) : (diameter >> 1); + const int radius = (diameter <= 0) ? static_cast(std::roundf(sigmaSpace * 1.5f)) : (diameter >> 1); float spaceCoeff = -1 / (2 * sigmaSpace * sigmaSpace); float colorCoeff = -1 / (2 * sigmaColor * sigmaColor); @@ -89,9 +88,9 @@ void dispatch_bilateral_filter_border_mode(hipStream_t stream, const Tensor &inp for (int j = 0; j < divisor; j++) { for (int i = 0; i < dividend; i++) { - threads.push_back(std::thread(Kernels::Host::bilateral_filter, ImageWrapper>, - inputWrapper, outputWrapper, radius, rollingHeight, rollingWidth, - prevHeight, prevWidth, spaceCoeff, colorCoeff)); + threads.push_back(std::thread( + Kernels::Host::bilateral_filter>, inputWrapper, + outputWrapper, radius, rollingHeight, rollingWidth, prevHeight, prevWidth, spaceCoeff, colorCoeff)); prevWidth = rollingWidth; rollingWidth += factorW; } diff --git a/src/op_bnd_box.cpp b/src/op_bnd_box.cpp index 3b2443e0..abaff32d 100644 --- a/src/op_bnd_box.cpp +++ b/src/op_bnd_box.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/detail/hip_utils.hpp" #include "core/tensor.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/bnd_box_device.hpp" #include "kernels/host/bnd_box_host.hpp" @@ -44,8 +44,8 @@ BndBox::~BndBox() {} template void dispatch_bnd_box_dtype(hipStream_t stream, const Tensor &input, const Tensor &output, std::shared_ptr> rects, eDeviceType device) { - ImageWrapper inputWrapper(input); - ImageWrapper outputWrapper(output); + TensorWrapper inputWrapper(input); + TensorWrapper outputWrapper(output); auto width = inputWrapper.width(); auto height = inputWrapper.height(); diff --git a/src/op_brightness_contrast.cpp b/src/op_brightness_contrast.cpp index 8673ee05..2708adb9 100644 --- a/src/op_brightness_contrast.cpp +++ b/src/op_brightness_contrast.cpp @@ -28,7 +28,7 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/brightness_contrast_device.hpp" #include "kernels/host/brightness_contrast_host.hpp" @@ -88,8 +88,8 @@ void dispatch_brightness_contrast_channels(hipStream_t stream, const Tensor &inp using SRC_DT_NC = detail::MakeType; using DST_DT_NC = detail::MakeType; - ImageWrapper inputWrapper(input); - ImageWrapper outputWrapper(output); + TensorWrapper inputWrapper(input); + TensorWrapper outputWrapper(output); // Launch CPU/GPU kernel depending on requested device type. switch (device) { diff --git a/src/op_composite.cpp b/src/op_composite.cpp index 650a7bf1..c73ac51f 100644 --- a/src/op_composite.cpp +++ b/src/op_composite.cpp @@ -24,7 +24,7 @@ #include #include "common/validation_helpers.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/composite_device.hpp" #include "kernels/host/composite_host.hpp" @@ -33,10 +33,10 @@ namespace roccv { template void dispatch_composite_masktype(hipStream_t stream, const Tensor& foreground, const Tensor& background, const Tensor& mask, const Tensor& output, eDeviceType device) { - ImageWrapper fgWrapper(foreground); - ImageWrapper bgWrapper(background); - ImageWrapper maskWrapper(mask); - ImageWrapper outputWrapper(output); + TensorWrapper fgWrapper(foreground); + TensorWrapper bgWrapper(background); + TensorWrapper maskWrapper(mask); + TensorWrapper outputWrapper(output); switch (device) { case eDeviceType::GPU: { diff --git a/src/op_convert_to.cpp b/src/op_convert_to.cpp index f814498e..4a25fbfd 100644 --- a/src/op_convert_to.cpp +++ b/src/op_convert_to.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. #include #include -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "common/validation_helpers.hpp" #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" @@ -40,8 +40,8 @@ void dispatch_convert_to_channels(hipStream_t stream, const Tensor &input, const using SRC_DT_NC = detail::MakeType; using DST_DT_NC = detail::MakeType; - ImageWrapper inputWrapper(input); - ImageWrapper outputWrapper(output); + TensorWrapper inputWrapper(input); + TensorWrapper outputWrapper(output); using SRC_BT = detail::BaseType; using DST_BT = detail::BaseType; diff --git a/src/op_copy_make_border.cpp b/src/op_copy_make_border.cpp index feacfbd9..ffc3e0a7 100644 --- a/src/op_copy_make_border.cpp +++ b/src/op_copy_make_border.cpp @@ -25,7 +25,7 @@ #include "common/validation_helpers.hpp" #include "core/wrappers/border_wrapper.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "core/wrappers/interpolation_wrapper.hpp" #include "kernels/device/copy_make_border_device.hpp" #include "kernels/host/copy_make_border_host.hpp" @@ -38,8 +38,8 @@ namespace roccv { template void dispatch_copy_make_border_border_mode(hipStream_t stream, const Tensor& input, const Tensor& output, int32_t top, int32_t left, T border_value, eDeviceType device) { - BorderWrapper in_desc(input, border_value); - ImageWrapper out_desc(output); + auto in_desc = MakeBorderWrapper(TensorWrapper(input), border_value); + TensorWrapper out_desc(output); switch (device) { case eDeviceType::GPU: { @@ -83,8 +83,7 @@ void dispatch_copy_make_border(hipStream_t stream, const Tensor& input, const Te } void CopyMakeBorder::operator()(hipStream_t stream, const Tensor& input, const Tensor& output, int32_t top, - int32_t left, eBorderType border_mode, float4 border_value, - eDeviceType device) const { + int32_t left, eBorderType border_mode, float4 border_value, eDeviceType device) const { CHECK_TENSOR_DEVICE(input, device); CHECK_TENSOR_LAYOUT(input, eTensorLayout::TENSOR_LAYOUT_NHWC, eTensorLayout::TENSOR_LAYOUT_HWC); CHECK_TENSOR_DATATYPES(input, eDataType::DATA_TYPE_U8, eDataType::DATA_TYPE_S8, eDataType::DATA_TYPE_U16, diff --git a/src/op_custom_crop.cpp b/src/op_custom_crop.cpp index 3d1a7056..13e5017c 100644 --- a/src/op_custom_crop.cpp +++ b/src/op_custom_crop.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include #include "common/validation_helpers.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/custom_crop_device.hpp" #include "kernels/host/custom_crop_host.hpp" @@ -35,8 +35,8 @@ namespace roccv { template void dispatch_custom_crop_dtype(hipStream_t stream, const Tensor& input, const Tensor& output, Box_t cropRect, eDeviceType device) { - ImageWrapper inputWrapper(input); - ImageWrapper outputWrapper(output); + TensorWrapper inputWrapper(input); + TensorWrapper outputWrapper(output); switch (device) { case eDeviceType::GPU: { diff --git a/src/op_cvt_color.cpp b/src/op_cvt_color.cpp index 31abd103..7785b465 100644 --- a/src/op_cvt_color.cpp +++ b/src/op_cvt_color.cpp @@ -27,7 +27,7 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/tensor.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/cvt_color_device.hpp" #include "kernels/host/cvt_color_host.hpp" @@ -87,38 +87,38 @@ void CvtColor::operator()(hipStream_t stream, const Tensor &input, Tensor &outpu switch (conversionCode) { case eColorConversionCode::COLOR_BGR2GRAY: Kernels::Device::rgb_or_bgr_to_grayscale - <<>>(ImageWrapper(input), ImageWrapper(output)); + <<>>(TensorWrapper(input), TensorWrapper(output)); break; case eColorConversionCode::COLOR_RGB2GRAY: Kernels::Device::rgb_or_bgr_to_grayscale - <<>>(ImageWrapper(input), ImageWrapper(output)); + <<>>(TensorWrapper(input), TensorWrapper(output)); break; case eColorConversionCode::COLOR_BGR2RGB: case eColorConversionCode::COLOR_RGB2BGR: Kernels::Device::reorder - <<>>(ImageWrapper(input), ImageWrapper(output)); + <<>>(TensorWrapper(input), TensorWrapper(output)); break; case eColorConversionCode::COLOR_BGR2YUV: Kernels::Device::rgb_or_bgr_to_yuv<<>>( - ImageWrapper(input), ImageWrapper(output), 128.0f); + TensorWrapper(input), TensorWrapper(output), 128.0f); break; case eColorConversionCode::COLOR_RGB2YUV: Kernels::Device::rgb_or_bgr_to_yuv<<>>( - ImageWrapper(input), ImageWrapper(output), 128.0f); + TensorWrapper(input), TensorWrapper(output), 128.0f); break; case eColorConversionCode::COLOR_YUV2BGR: Kernels::Device::yuv_to_rgb_or_bgr<<>>( - ImageWrapper(input), ImageWrapper(output), 128.0f); + TensorWrapper(input), TensorWrapper(output), 128.0f); break; case eColorConversionCode::COLOR_YUV2RGB: Kernels::Device::yuv_to_rgb_or_bgr<<>>( - ImageWrapper(input), ImageWrapper(output), 128.0f); + TensorWrapper(input), TensorWrapper(output), 128.0f); break; default: @@ -129,39 +129,39 @@ void CvtColor::operator()(hipStream_t stream, const Tensor &input, Tensor &outpu switch (conversionCode) { case eColorConversionCode::COLOR_BGR2GRAY: - Kernels::Host::rgb_or_bgr_to_grayscale(ImageWrapper(input), - ImageWrapper(output)); + Kernels::Host::rgb_or_bgr_to_grayscale(TensorWrapper(input), + TensorWrapper(output)); break; case eColorConversionCode::COLOR_RGB2GRAY: - Kernels::Host::rgb_or_bgr_to_grayscale(ImageWrapper(input), - ImageWrapper(output)); + Kernels::Host::rgb_or_bgr_to_grayscale(TensorWrapper(input), + TensorWrapper(output)); break; case eColorConversionCode::COLOR_BGR2RGB: case eColorConversionCode::COLOR_RGB2BGR: - Kernels::Host::reorder(ImageWrapper(input), - ImageWrapper(output)); + Kernels::Host::reorder(TensorWrapper(input), + TensorWrapper(output)); break; case eColorConversionCode::COLOR_BGR2YUV: - Kernels::Host::rgb_or_bgr_to_yuv(ImageWrapper(input), - ImageWrapper(output), 128.0f); + Kernels::Host::rgb_or_bgr_to_yuv(TensorWrapper(input), + TensorWrapper(output), 128.0f); break; case eColorConversionCode::COLOR_RGB2YUV: - Kernels::Host::rgb_or_bgr_to_yuv(ImageWrapper(input), - ImageWrapper(output), 128.0f); + Kernels::Host::rgb_or_bgr_to_yuv(TensorWrapper(input), + TensorWrapper(output), 128.0f); break; case eColorConversionCode::COLOR_YUV2BGR: - Kernels::Host::yuv_to_rgb_or_bgr(ImageWrapper(input), - ImageWrapper(output), 128.0f); + Kernels::Host::yuv_to_rgb_or_bgr(TensorWrapper(input), + TensorWrapper(output), 128.0f); break; case eColorConversionCode::COLOR_YUV2RGB: - Kernels::Host::yuv_to_rgb_or_bgr(ImageWrapper(input), - ImageWrapper(output), 128.0f); + Kernels::Host::yuv_to_rgb_or_bgr(TensorWrapper(input), + TensorWrapper(output), 128.0f); break; default: diff --git a/src/op_flip.cpp b/src/op_flip.cpp index 5566a0b6..d28ab8fa 100644 --- a/src/op_flip.cpp +++ b/src/op_flip.cpp @@ -29,7 +29,7 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/exception.hpp" #include "core/status_type.h" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/flip_device.hpp" #include "kernels/host/flip_host.hpp" @@ -37,8 +37,8 @@ namespace roccv { template void dispatch_flip_axis(hipStream_t stream, const Tensor& input, const Tensor& output, eDeviceType device) { - ImageWrapper inputWrapper(input); - ImageWrapper outputWrapper(output); + TensorWrapper inputWrapper(input); + TensorWrapper outputWrapper(output); switch (device) { case eDeviceType::GPU: { diff --git a/src/op_gamma_contrast.cpp b/src/op_gamma_contrast.cpp index d6c78690..e4ca3520 100644 --- a/src/op_gamma_contrast.cpp +++ b/src/op_gamma_contrast.cpp @@ -34,7 +34,7 @@ THE SOFTWARE. #include "common/math_vector.hpp" #include "common/validation_helpers.hpp" #include "core/tensor.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/gamma_contrast_device.hpp" #include "kernels/host/gamma_contrast_host.hpp" @@ -43,8 +43,8 @@ namespace roccv { template void dispatch_gamma_contrast_dtype(hipStream_t stream, const Tensor &input, const Tensor &output, float gamma, eDeviceType device) { - ImageWrapper inputWrapper(input); - ImageWrapper outputWrapper(output); + TensorWrapper inputWrapper(input); + TensorWrapper outputWrapper(output); if (device == eDeviceType::GPU) { dim3 block(64, 16); diff --git a/src/op_histogram.cpp b/src/op_histogram.cpp index 24a2ef79..0627a732 100644 --- a/src/op_histogram.cpp +++ b/src/op_histogram.cpp @@ -31,7 +31,7 @@ THE SOFTWARE. #include "common/array_wrapper.hpp" #include "common/validation_helpers.hpp" #include "core/wrappers/generic_tensor_wrapper.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/histogram_device.hpp" #include "kernels/host/histogram_host.hpp" @@ -44,7 +44,7 @@ template void dispatch_histogram_dtype(hipStream_t stream, const Tensor& input, std::optional> mask, const Tensor& histogram, eDeviceType device) { - ImageWrapper inputWrapper(input); + TensorWrapper inputWrapper(input); const auto o_height = histogram.shape()[histogram.shape().layout().height_index()]; const auto o_width = histogram.shape()[histogram.shape().layout().width_index()]; @@ -91,7 +91,7 @@ void dispatch_histogram_dtype(hipStream_t stream, const Tensor& input, std::reference_wrapper mask_ref = mask.value(); const Tensor& actual_mask = mask_ref.get(); CHECK_TENSOR_COMPARISON(input.shape() == actual_mask.shape()); - ImageWrapper maskWrapper(actual_mask); + TensorWrapper maskWrapper(actual_mask); Kernels::Device::histogram_kernel<<>>( inputWrapper, maskWrapper, GenericTensorWrapper(histogram)); } else { @@ -115,7 +115,7 @@ void dispatch_histogram_dtype(hipStream_t stream, const Tensor& input, std::reference_wrapper mask_ref = mask.value(); const Tensor& actual_mask = mask_ref.get(); CHECK_TENSOR_COMPARISON(input.shape() == actual_mask.shape()); - ImageWrapper maskWrapper(actual_mask); + TensorWrapper maskWrapper(actual_mask); Kernels::Host::histogram_kernel(inputWrapper, maskWrapper, GenericTensorWrapper(histogram)); } else { Kernels::Host::histogram_kernel(inputWrapper, GenericTensorWrapper(histogram)); diff --git a/src/op_normalize.cpp b/src/op_normalize.cpp index 529638b8..ca2fde96 100644 --- a/src/op_normalize.cpp +++ b/src/op_normalize.cpp @@ -28,7 +28,7 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/detail/type_traits.hpp" #include "core/tensor.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/normalize_device.hpp" #include "kernels/host/normalize_host.hpp" @@ -42,10 +42,10 @@ void dispatch_normalize_stddev(hipStream_t stream, const Tensor& input, const Te // tensors. using work_type = detail::MakeType>; - ImageWrapper inputWrap(input); - ImageWrapper outputWrap(output); - ImageWrapper scaleWrap(scale); - ImageWrapper baseWrap(base); + TensorWrapper inputWrap(input); + TensorWrapper outputWrap(output); + TensorWrapper scaleWrap(scale); + TensorWrapper baseWrap(base); switch (device) { case eDeviceType::GPU: { diff --git a/src/op_reformat.cpp b/src/op_reformat.cpp index 4a79145b..cda3c809 100644 --- a/src/op_reformat.cpp +++ b/src/op_reformat.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include #include "common/validation_helpers.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "kernels/device/reformat_device.hpp" #include "kernels/host/reformat_host.hpp" @@ -35,8 +35,8 @@ namespace roccv { namespace { template void DispatchReformatChannels(hipStream_t stream, const Tensor& input, const Tensor& output, eDeviceType device) { - ImageWrapper inputWrap(input); - ImageWrapper outputWrap(output); + TensorWrapper inputWrap(input); + TensorWrapper outputWrap(output); switch (device) { case eDeviceType::GPU: { diff --git a/src/op_remap.cpp b/src/op_remap.cpp index 0992cf44..40d14427 100644 --- a/src/op_remap.cpp +++ b/src/op_remap.cpp @@ -29,7 +29,7 @@ THE SOFTWARE. #include "core/detail/internal_structs.hpp" #include "core/detail/math/math.hpp" #include "core/detail/type_traits.hpp" -#include "core/wrappers/image_wrapper.hpp" +#include "core/wrappers/tensor_wrapper.hpp" #include "core/wrappers/interpolation_wrapper.hpp" #include "kernels/device/remap_device.hpp" #include "kernels/host/remap_host.hpp" @@ -76,9 +76,10 @@ template void dispatch_remap_mapInterp(hipStream_t stream, const Tensor &input, const Tensor &output, const Tensor &map, const eRemapType mapValueType, const bool alignCorners, const T borderValue, const eDeviceType device) { - ImageWrapper outputWrapper(output); - InterpolationWrapper wrappedMapTensor(map, make_float2(0, 0)); - InterpolationWrapper inputWrapper(input, borderValue); + TensorWrapper outputWrapper(output); + auto wrappedMapTensor = + MakeInterpolationWrapper(MakeBorderWrapper(TensorWrapper(map), make_float2(0, 0))); + auto inputWrapper = MakeInterpolationWrapper(MakeBorderWrapper(TensorWrapper(input), borderValue)); int mapBatchSize = wrappedMapTensor.batches(); diff --git a/src/op_resize.cpp b/src/op_resize.cpp index d7cd0b61..dda08eaf 100644 --- a/src/op_resize.cpp +++ b/src/op_resize.cpp @@ -25,50 +25,88 @@ THE SOFTWARE. #include #include "common/validation_helpers.hpp" -#include "core/detail/casting.hpp" #include "core/exception.hpp" +#include "core/image_batch_var_shape.hpp" #include "core/status_type.h" +#include "core/wrappers/image_batch_var_shape_wrapper.hpp" #include "core/wrappers/interpolation_wrapper.hpp" #include "kernels/device/resize_device.hpp" #include "kernels/host/resize_host.hpp" namespace roccv { -template -void dispatch_resize_interp(hipStream_t stream, const Tensor& input, const Tensor& output, eDeviceType device) { - ImageWrapper outputWrapper(output); +// Common launch tail shared by the uniform (Tensor) and variable-shape (ImageBatchVarShape) input paths. The +// source base wrapper (TensorWrapper or ImageBatchVarShapeWrapper) is built by the caller; everything +// downstream — border/interpolation composition, output wrapper, grid, kernel — is identical. The kernel derives +// the per-image scale internally, so no scale is computed here. +template +void dispatch_resize_launch(hipStream_t stream, SrcBaseWrapper srcBase, const Tensor& output, eDeviceType device) { + TensorWrapper outputWrapper(output); // Resize operation should clamp values at the border (REPLICATE border mode) - InterpolationWrapper inputWrapper(input, T{}); - - float scaleX = inputWrapper.width() / static_cast(outputWrapper.width()); - float scaleY = inputWrapper.height() / static_cast(outputWrapper.height()); + auto inputWrapper = + MakeInterpolationWrapper(MakeBorderWrapper(srcBase, T{})); switch (device) { case eDeviceType::GPU: { dim3 block(64, 16); dim3 grid((outputWrapper.width() + block.x - 1) / block.x, (outputWrapper.height() + block.y - 1) / block.y, outputWrapper.batches()); - Kernels::Device::resize<<>>(inputWrapper, outputWrapper, scaleX, scaleY); + Kernels::Device::resize<<>>(inputWrapper, outputWrapper); break; } case eDeviceType::CPU: { - Kernels::Host::resize(inputWrapper, outputWrapper, scaleX, scaleY); + Kernels::Host::resize(inputWrapper, outputWrapper); break; } } } +template +void dispatch_resize_interp(hipStream_t stream, const Tensor& input, const Tensor& output, eDeviceType device) { + dispatch_resize_launch(stream, TensorWrapper(input), output, device); +} + +template +void dispatch_resize_interp_var(hipStream_t stream, ImageBatchVarShape& input, const Tensor& output, + eDeviceType device) { + // The exported snapshot is a non-owning view into the batch's descriptor table; the wrapper copies out the + // device/host pointer it needs, and the batch outlives this call. + auto data = input.exportData(stream); + dispatch_resize_launch(stream, ImageBatchVarShapeWrapper(data), output, device); +} + template void dispatch_resize_dtype(hipStream_t stream, const Tensor& input, const Tensor& output, eInterpolationType interpolation, eDeviceType device) { + static const std::unordered_map> + funcs = { + {eInterpolationType::INTERP_TYPE_NEAREST, + dispatch_resize_interp}, + {eInterpolationType::INTERP_TYPE_LINEAR, dispatch_resize_interp}, + {eInterpolationType::INTERP_TYPE_CUBIC, dispatch_resize_interp}}; + + if (!funcs.contains(interpolation)) { + throw Exception("Operation does not support the given interpolation mode.", eStatusType::NOT_IMPLEMENTED); + } + + auto func = funcs.at(interpolation); + func(stream, input, output, device); +} + +template +void dispatch_resize_dtype_var(hipStream_t stream, ImageBatchVarShape& input, const Tensor& output, + eInterpolationType interpolation, eDeviceType device) { static const std::unordered_map< eInterpolationType, - std::function> - funcs = {{eInterpolationType::INTERP_TYPE_NEAREST, dispatch_resize_interp}, - {eInterpolationType::INTERP_TYPE_LINEAR, dispatch_resize_interp}, - {eInterpolationType::INTERP_TYPE_CUBIC, dispatch_resize_interp} - }; + std::function> + funcs = {{eInterpolationType::INTERP_TYPE_NEAREST, + dispatch_resize_interp_var}, + {eInterpolationType::INTERP_TYPE_LINEAR, + dispatch_resize_interp_var}, + {eInterpolationType::INTERP_TYPE_CUBIC, + dispatch_resize_interp_var}}; if (!funcs.contains(interpolation)) { throw Exception("Operation does not support the given interpolation mode.", eStatusType::NOT_IMPLEMENTED); @@ -78,8 +116,8 @@ void dispatch_resize_dtype(hipStream_t stream, const Tensor& input, const Tensor func(stream, input, output, device); } -void Resize::operator()(hipStream_t stream, const Tensor& input, const Tensor& output, - eInterpolationType interpolation, eDeviceType device) const { +void Resize::operator()(hipStream_t stream, const Tensor& input, const Tensor& output, eInterpolationType interpolation, + eDeviceType device) const { CHECK_TENSOR_DEVICE(input, device); CHECK_TENSOR_DEVICE(output, device); @@ -109,4 +147,40 @@ void Resize::operator()(hipStream_t stream, const Tensor& input, const Tensor& o if (func == 0) throw Exception("Not mapped to a defined function.", eStatusType::INVALID_OPERATION); func(stream, input, output, interpolation, device); } + +void Resize::operator()(hipStream_t stream, ImageBatchVarShape& input, const Tensor& output, + eInterpolationType interpolation, eDeviceType device) const { + // The variable-shape batch resolves dtype/channels from its single shared format; a heterogeneous (or empty) + // batch can't be expressed by the T-templated wrapper and is rejected. + CHECK_IMAGE_BATCH_DEVICE(input, device); + CHECK_TENSOR_DEVICE(output, device); + + CHECK_IMAGE_BATCH_UNIFORM_FORMAT(input); + CHECK_IMAGE_BATCH_DATATYPES(input, DATA_TYPE_U8, DATA_TYPE_F32); + CHECK_IMAGE_BATCH_CHANNELS(input, 1, 3, 4); + + // The output holds the resized batch as a uniform tensor, so it must carry an explicit batch dimension. + CHECK_TENSOR_LAYOUT(output, TENSOR_LAYOUT_NHWC); + CHECK_TENSOR_DATATYPES(output, DATA_TYPE_U8, DATA_TYPE_F32); + CHECK_TENSOR_CHANNELS(output, 1, 3, 4); + + // Output dtype/channels/batch must agree with the input batch. + ImageFormat format = input.uniqueFormat(); + CHECK_TENSOR_COMPARISON(output.dtype().etype() == format.dtype()); + CHECK_TENSOR_COMPARISON(output.shape(output.layout().channels_index()) == format.channels()); + CHECK_TENSOR_COMPARISON(output.shape(output.layout().batch_index()) == input.numImages()); + + // clang-format off + static const std::unordered_map, 4>> + funcs = { + {eDataType::DATA_TYPE_U8, {dispatch_resize_dtype_var, 0, dispatch_resize_dtype_var, dispatch_resize_dtype_var}}, + {eDataType::DATA_TYPE_F32, {dispatch_resize_dtype_var, 0, dispatch_resize_dtype_var, dispatch_resize_dtype_var}} + }; + // clang-format on + + auto func = funcs.at(format.dtype())[format.channels() - 1]; + if (func == 0) throw Exception("Not mapped to a defined function.", eStatusType::INVALID_OPERATION); + func(stream, input, output, interpolation, device); +} } // namespace roccv diff --git a/src/op_rotate.cpp b/src/op_rotate.cpp index 28806779..508de16c 100644 --- a/src/op_rotate.cpp +++ b/src/op_rotate.cpp @@ -54,8 +54,9 @@ void dispatch_rotate_interp(hipStream_t stream, const Tensor &input, const Tenso T borderVal = detail::SaturateCast(make_float4(0.0f, 0.0f, 0.0f, 0.0f)); - ImageWrapper outputWrap(output); - InterpolationWrapper inputWrap(input, borderVal); + TensorWrapper outputWrap(output); + auto inputWrap = MakeInterpolationWrapper( + MakeBorderWrapper(TensorWrapper(input), borderVal)); switch (device) { case eDeviceType::GPU: { @@ -74,8 +75,8 @@ void dispatch_rotate_interp(hipStream_t stream, const Tensor &input, const Tenso } template -void dispatch_rotate_type(hipStream_t stream, const Tensor &input, const Tensor &output, double angleDeg, - double2 shift, eInterpolationType interpolation, eDeviceType device) { +void dispatch_rotate_type(hipStream_t stream, const Tensor &input, const Tensor &output, double angleDeg, double2 shift, + eInterpolationType interpolation, eDeviceType device) { // clang-format off static const std::unordered_map void dispatch_threshold_dtype(hipStream_t stream, const Tensor &input, const Tensor &output, const Tensor &thresh, const Tensor &maxVal, eThresholdType m_threshType, eDeviceType device) { - ImageWrapper inputWrapper(input); - ImageWrapper outputWrapper(output); + TensorWrapper inputWrapper(input); + TensorWrapper outputWrapper(output); const auto height = input.shape()[input.shape().layout().height_index()]; const auto width = input.shape()[input.shape().layout().width_index()]; diff --git a/src/op_warp_affine.cpp b/src/op_warp_affine.cpp index ee29111f..df3a420d 100644 --- a/src/op_warp_affine.cpp +++ b/src/op_warp_affine.cpp @@ -35,8 +35,8 @@ template void dispatch_warp_affine_interp(hipStream_t stream, const Tensor &input, const Tensor &output, const AffineTransform affineInv, T borderValue, eDeviceType device) { ArrayWrapper transform(affineInv); - ImageWrapper outputWrapper(output); - InterpolationWrapper inputWrapper(input, borderValue); + TensorWrapper outputWrapper(output); + auto inputWrapper = MakeInterpolationWrapper(MakeBorderWrapper(TensorWrapper(input), borderValue)); switch (device) { case eDeviceType::GPU: { diff --git a/src/op_warp_perspective.cpp b/src/op_warp_perspective.cpp index ca77fc8c..1a7e12f2 100644 --- a/src/op_warp_perspective.cpp +++ b/src/op_warp_perspective.cpp @@ -36,8 +36,8 @@ template void dispatch_warp_perspective_interp(hipStream_t stream, const Tensor &input, const Tensor &output, const PerspectiveTransform transMatrix, T borderValue, eDeviceType device) { ArrayWrapper transform(transMatrix); - ImageWrapper outputWrapper(output); - InterpolationWrapper inputWrapper(input, borderValue); + TensorWrapper outputWrapper(output); + auto inputWrapper = MakeInterpolationWrapper(MakeBorderWrapper(TensorWrapper(input), borderValue)); // Launch CPU/GPU kernel depending on requested device type. switch (device) { diff --git a/tests/roccv/cpp/src/tests/core/wrappers/test_border_wrapper.cpp b/tests/roccv/cpp/src/tests/core/wrappers/test_border_wrapper.cpp index 873f05dc..6dd736ad 100644 --- a/tests/roccv/cpp/src/tests/core/wrappers/test_border_wrapper.cpp +++ b/tests/roccv/cpp/src/tests/core/wrappers/test_border_wrapper.cpp @@ -103,7 +103,7 @@ int64_t GetCoordOfBorderPel(int64_t u, int64_t dimSize, eBorderType borderMode) * * @tparam T The underlying datatype of the image. (e.g. uchar3) * @tparam BT The base datatype of the image (e.g. unsigned char) - * @param[in] input The input ImageWrapper referencing the underlying image data. + * @param[in] input The input TensorWrapper referencing the underlying image data. * @param[in] borderMode The border mode used to handle out of bounds coordinates. * @param[in] borderValue The value to fallback to when handling out of bounds coordinates with the CONSTANT border * mode. @@ -115,8 +115,8 @@ int64_t GetCoordOfBorderPel(int64_t u, int64_t dimSize, eBorderType borderMode) * coordinates fall out of bounds. */ template > -BT GoldenBorderAt(ImageWrapper& input, eBorderType borderMode, T borderValue, int64_t sample, int64_t y, - int64_t x, int64_t channel) { +BT GoldenBorderAt(TensorWrapper& input, eBorderType borderMode, T borderValue, int64_t sample, int64_t y, int64_t x, + int64_t channel) { int64_t outX = x, outY = y; if (borderMode == eBorderType::BORDER_TYPE_CONSTANT) { @@ -130,7 +130,7 @@ BT GoldenBorderAt(ImageWrapper& input, eBorderType borderMode, T borderValue, outY = GetCoordOfBorderPel(y, input.height(), borderMode); } - // Return the value at the modified outX, outY coordinates using the passed in ImageWrapper. + // Return the value at the modified outX, outY coordinates using the passed in TensorWrapper. return detail::GetElement(input.at(sample, outY, outX, 0), channel); } @@ -161,7 +161,8 @@ void TestCorrectness(float4 borderValue, int32_t batchSize, Size2D imageSize, in FillVector(inputData); // BorderWrapper to calculate the actual calculated values. - BorderWrapper borderWrap(ImageWrapper(inputData, batchSize, imageSize.w, imageSize.h), borderVal); + auto borderWrap = + MakeBorderWrapper(TensorWrapper(inputData, batchSize, imageSize.w, imageSize.h), borderVal); std::vector actualOutput(numElementsWithBorder); int actualIndex = 0; for (int batch = 0; batch < batchSize; ++batch) { @@ -176,9 +177,9 @@ void TestCorrectness(float4 borderValue, int32_t batchSize, Size2D imageSize, in } } - // ImageWrapper for use in the golden output generator. ImageWrapper is unit tested separately, and is + // TensorWrapper for use in the golden output generator. TensorWrapper is unit tested separately, and is // considered working at this point in the dependency chain. - ImageWrapper imageWrap(inputData, batchSize, imageSize.w, imageSize.h); + TensorWrapper imageWrap(inputData, batchSize, imageSize.w, imageSize.h); std::vector goldenOutput(numElementsWithBorder); int goldenIndex = 0; for (int batch = 0; batch < batchSize; ++batch) { diff --git a/tests/roccv/cpp/src/tests/core/wrappers/test_image_batch_var_shape_wrapper.cpp b/tests/roccv/cpp/src/tests/core/wrappers/test_image_batch_var_shape_wrapper.cpp new file mode 100644 index 00000000..12d5000e --- /dev/null +++ b/tests/roccv/cpp/src/tests/core/wrappers/test_image_batch_var_shape_wrapper.cpp @@ -0,0 +1,386 @@ +/* + * 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 +#include +#include +#include + +#include "test_helpers.hpp" + +using namespace roccv; +using namespace roccv::tests; + +namespace { + +// Per-image copy kernel: writes dst[n,y,x,c] = src[n,y,x,c] for the image at batch index n. +// Each launch covers exactly one image — the host loop steps through n and resizes the +// grid to that image's dimensions, which avoids needing a max-bounds check inside the kernel. +template +__global__ void VarShapeCopyKernel(ImageBatchVarShapeWrapper src, ImageBatchVarShapeWrapper dst, int32_t n) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst.width(n) || y >= dst.height(n)) return; + dst.at(n, y, x, 0) = src.at(n, y, x, 0); +} + +// Roundtrip test: build a varshape batch from heterogeneous host pixel data, run the copy kernel from src varshape +// wrapper to dst varshape wrapper, read pixels back, and verify byte-equality with the original input. +template > +void TestRoundtripCopy(const std::vector& sizes, ImageFormat fmt) { + const int channels = detail::NumElements; + const int32_t numImages = static_cast(sizes.size()); + + // Generate per-image host pixel data. + std::vector> hostPixels(numImages); + for (int32_t i = 0; i < numImages; ++i) { + hostPixels[i].resize(static_cast(sizes[i].w) * sizes[i].h * channels); + FillVector(hostPixels[i], /*seed=*/static_cast(0x1000 + i)); + } + + hipStream_t stream = nullptr; + + // Build source batch and copy host pixels in. + ImageBatchVarShape srcBatch(numImages); + std::vector srcImages; + srcImages.reserve(numImages); + for (int32_t i = 0; i < numImages; ++i) { + srcImages.emplace_back(sizes[i], fmt); + auto sd = srcImages[i].exportData(); + const ImagePlaneStrided& sp = sd.plane(0); + const size_t rowBytes = static_cast(sizes[i].w) * channels * sizeof(BT); + HIP_VALIDATE_NO_ERRORS(hipMemcpy2DAsync(sp.basePtr, sp.rowStride, hostPixels[i].data(), rowBytes, rowBytes, + sizes[i].h, hipMemcpyHostToDevice, stream)); + srcBatch.pushBack(srcImages[i]); + } + + // Build destination batch with matching shapes. + ImageBatchVarShape dstBatch(numImages); + std::vector dstImages; + dstImages.reserve(numImages); + for (int32_t i = 0; i < numImages; ++i) { + dstImages.emplace_back(sizes[i], fmt); + dstBatch.pushBack(dstImages[i]); + } + + auto srcData = srcBatch.exportData(stream); + auto dstData = dstBatch.exportData(stream); + ImageBatchVarShapeWrapper srcWrap(srcData); + ImageBatchVarShapeWrapper dstWrap(dstData); + + // Launch one kernel per image (sizes vary so a single 3D launch can't bound y to per-image height cleanly). + for (int32_t i = 0; i < numImages; ++i) { + dim3 block(16, 16); + dim3 grid((sizes[i].w + block.x - 1) / block.x, (sizes[i].h + block.y - 1) / block.y); + VarShapeCopyKernel<<>>(srcWrap, dstWrap, i); + } + + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream)); + + // Read back dst pixels and verify byte-for-byte against the original host input. + for (int32_t i = 0; i < numImages; ++i) { + std::vector dstHost(static_cast(sizes[i].w) * sizes[i].h * channels); + auto dd = dstImages[i].exportData(); + const ImagePlaneStrided& dp = dd.plane(0); + const size_t rowBytes = static_cast(sizes[i].w) * channels * sizeof(BT); + HIP_VALIDATE_NO_ERRORS(hipMemcpy2D(dstHost.data(), rowBytes, dp.basePtr, dp.rowStride, rowBytes, sizes[i].h, + hipMemcpyDeviceToHost)); + CompareVectors(dstHost, hostPixels[i]); + } +} + +// Border-composition test: write the BORDER_TYPE_CONSTANT fallback for every output pixel by reading from a coordinate +// that is guaranteed to be out of bounds for every image (-1, -1). Confirms BorderWrapper> correctly delegates to width(n) / height(n) for per-image bounds; otherwise it would +// dereference invalid memory. +template +__global__ void VarShapeBorderConstantKernel( + BorderWrapper> src, + ImageBatchVarShapeWrapper dst, int32_t n) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst.width(n) || y >= dst.height(n)) return; + dst.at(n, y, x, 0) = src.at(n, -1, -1, 0); +} + +template > +void TestBorderConstantComposition(const std::vector& sizes, ImageFormat fmt, T borderValue) { + const int channels = detail::NumElements; + const int32_t numImages = static_cast(sizes.size()); + + // Source pixels content doesn't matter — every read is forced OOB. + ImageBatchVarShape srcBatch(numImages); + ImageBatchVarShape dstBatch(numImages); + std::vector srcImages, dstImages; + srcImages.reserve(numImages); + dstImages.reserve(numImages); + for (int32_t i = 0; i < numImages; ++i) { + srcImages.emplace_back(sizes[i], fmt); + dstImages.emplace_back(sizes[i], fmt); + srcBatch.pushBack(srcImages[i]); + dstBatch.pushBack(dstImages[i]); + } + + hipStream_t stream = nullptr; + auto srcData = srcBatch.exportData(stream); + auto dstData = dstBatch.exportData(stream); + + auto srcWrap = + MakeBorderWrapper(ImageBatchVarShapeWrapper(srcData), borderValue); + ImageBatchVarShapeWrapper dstWrap(dstData); + + for (int32_t i = 0; i < numImages; ++i) { + dim3 block(16, 16); + dim3 grid((sizes[i].w + block.x - 1) / block.x, (sizes[i].h + block.y - 1) / block.y); + VarShapeBorderConstantKernel<<>>(srcWrap, dstWrap, i); + } + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream)); + + // Expect every output pixel of every image to equal borderValue. + std::vector borderBytes(channels); + for (int c = 0; c < channels; ++c) borderBytes[c] = detail::GetElement(borderValue, c); + + for (int32_t i = 0; i < numImages; ++i) { + const size_t pixels = static_cast(sizes[i].w) * sizes[i].h; + std::vector dstHost(pixels * channels); + auto dd = dstImages[i].exportData(); + const ImagePlaneStrided& dp = dd.plane(0); + const size_t rowBytes = static_cast(sizes[i].w) * channels * sizeof(BT); + HIP_VALIDATE_NO_ERRORS(hipMemcpy2D(dstHost.data(), rowBytes, dp.basePtr, dp.rowStride, rowBytes, sizes[i].h, + hipMemcpyDeviceToHost)); + std::vector expected(pixels * channels); + for (size_t p = 0; p < pixels; ++p) { + for (int c = 0; c < channels; ++c) expected[p * channels + c] = borderBytes[c]; + } + CompareVectors(dstHost, expected); + } +} + +// Interpolation-composition test: NEAREST interpolation at integer coordinates is the identity, so a roundtrip copy +// via InterpolationWrapper must equal the source. Confirms the full wrapper chain +// composes correctly over a VarShape backing. +template +__global__ void VarShapeInterpNearestKernel( + InterpolationWrapper>> + src, + ImageBatchVarShapeWrapper dst, int32_t n) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst.width(n) || y >= dst.height(n)) return; + dst.at(n, y, x, 0) = src.at(n, static_cast(y), static_cast(x), 0); +} + +template > +void TestInterpolationNearestComposition(const std::vector& sizes, ImageFormat fmt) { + const int channels = detail::NumElements; + const int32_t numImages = static_cast(sizes.size()); + + std::vector> hostPixels(numImages); + for (int32_t i = 0; i < numImages; ++i) { + hostPixels[i].resize(static_cast(sizes[i].w) * sizes[i].h * channels); + FillVector(hostPixels[i], static_cast(0x2000 + i)); + } + + hipStream_t stream = nullptr; + ImageBatchVarShape srcBatch(numImages); + ImageBatchVarShape dstBatch(numImages); + std::vector srcImages, dstImages; + srcImages.reserve(numImages); + dstImages.reserve(numImages); + for (int32_t i = 0; i < numImages; ++i) { + srcImages.emplace_back(sizes[i], fmt); + auto sd = srcImages[i].exportData(); + const ImagePlaneStrided& sp = sd.plane(0); + const size_t rowBytes = static_cast(sizes[i].w) * channels * sizeof(BT); + HIP_VALIDATE_NO_ERRORS(hipMemcpy2DAsync(sp.basePtr, sp.rowStride, hostPixels[i].data(), rowBytes, rowBytes, + sizes[i].h, hipMemcpyHostToDevice, stream)); + srcBatch.pushBack(srcImages[i]); + + dstImages.emplace_back(sizes[i], fmt); + dstBatch.pushBack(dstImages[i]); + } + + auto srcData = srcBatch.exportData(stream); + auto dstData = dstBatch.exportData(stream); + + auto srcWrap = MakeInterpolationWrapper( + MakeBorderWrapper(ImageBatchVarShapeWrapper(srcData), T{})); + ImageBatchVarShapeWrapper dstWrap(dstData); + + for (int32_t i = 0; i < numImages; ++i) { + dim3 block(16, 16); + dim3 grid((sizes[i].w + block.x - 1) / block.x, (sizes[i].h + block.y - 1) / block.y); + VarShapeInterpNearestKernel<<>>(srcWrap, dstWrap, i); + } + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream)); + + for (int32_t i = 0; i < numImages; ++i) { + std::vector dstHost(static_cast(sizes[i].w) * sizes[i].h * channels); + auto dd = dstImages[i].exportData(); + const ImagePlaneStrided& dp = dd.plane(0); + const size_t rowBytes = static_cast(sizes[i].w) * channels * sizeof(BT); + HIP_VALIDATE_NO_ERRORS(hipMemcpy2D(dstHost.data(), rowBytes, dp.basePtr, dp.rowStride, rowBytes, sizes[i].h, + hipMemcpyDeviceToHost)); + CompareVectors(dstHost, hostPixels[i]); + } +} + +// Verify accessor surface: width(n), height(n), batches(), channels(). +template +void TestAccessors(const std::vector& sizes, ImageFormat fmt) { + const int32_t numImages = static_cast(sizes.size()); + ImageBatchVarShape batch(numImages); + std::vector handles; + handles.reserve(numImages); + for (int32_t i = 0; i < numImages; ++i) { + handles.emplace_back(sizes[i], fmt); + batch.pushBack(handles[i]); + } + auto data = batch.exportData(0); + ImageBatchVarShapeWrapper wrap(data); + + EXPECT_EQ(wrap.batches(), static_cast(numImages)); + EXPECT_EQ(wrap.channels(), static_cast(detail::NumElements)); + // width/height are device pointers under the hood; reading them on host post-sync is safe because exportData + // recorded a hipEvent that hipStreamSynchronize on the null stream above (implicit) drains. The descriptor table + // lives in device memory though, so we round-trip the lookups through a tiny D->H pull via the wrapper's host + // mirror path — here we just check that the construction succeeded; per-image width/height behavior is exercised + // end-to-end by TestRoundtripCopy. + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(0)); +} + +// Host-path roundtrip: a CPU-resident batch exports a host snapshot, and the wrapper reads/writes +// it directly from host code — no kernel launch, no stream. Mirrors TestRoundtripCopy for the GPU +// path and confirms the residency-agnostic constructor handles the ...Host leaf. +template > +void TestHostRoundtripCopy(const std::vector& sizes, ImageFormat fmt) { + const int channels = detail::NumElements; + const int32_t numImages = static_cast(sizes.size()); + + std::vector> hostPixels(numImages); + for (int32_t i = 0; i < numImages; ++i) { + hostPixels[i].resize(static_cast(sizes[i].w) * sizes[i].h * channels); + FillVector(hostPixels[i], /*seed=*/static_cast(0x2000 + i)); + } + + // CPU-resident batches; fill each source image's host buffer directly. + ImageBatchVarShape srcBatch(numImages, eDeviceType::CPU); + ImageBatchVarShape dstBatch(numImages, eDeviceType::CPU); + std::vector srcImages, dstImages; + srcImages.reserve(numImages); + dstImages.reserve(numImages); + for (int32_t i = 0; i < numImages; ++i) { + srcImages.emplace_back(sizes[i], fmt, eDeviceType::CPU); + dstImages.emplace_back(sizes[i], fmt, eDeviceType::CPU); + + auto sd = srcImages[i].exportData(); + const ImagePlaneStrided& sp = sd.plane(0); + const size_t rowBytes = static_cast(sizes[i].w) * channels * sizeof(BT); + for (int32_t y = 0; y < sizes[i].h; ++y) { + std::memcpy(static_cast(sp.basePtr) + y * sp.rowStride, + hostPixels[i].data() + static_cast(y) * sizes[i].w * channels, rowBytes); + } + + srcBatch.pushBack(srcImages[i]); + dstBatch.pushBack(dstImages[i]); + } + + auto srcData = srcBatch.exportData(0); + auto dstData = dstBatch.exportData(0); + ImageBatchVarShapeWrapper srcWrap(srcData); + ImageBatchVarShapeWrapper dstWrap(dstData); + + // Copy whole pixels through the wrapper entirely on the host (same at(n,y,x,0) semantics as + // VarShapeCopyKernel, just without a device launch). + for (int32_t n = 0; n < numImages; ++n) { + for (int64_t y = 0; y < srcWrap.height(n); ++y) { + for (int64_t x = 0; x < srcWrap.width(n); ++x) { + dstWrap.at(n, y, x, 0) = srcWrap.at(n, y, x, 0); + } + } + } + + // Read dst host buffers back and verify byte-for-byte against the original input. + for (int32_t i = 0; i < numImages; ++i) { + std::vector dstHost(static_cast(sizes[i].w) * sizes[i].h * channels); + auto dd = dstImages[i].exportData(); + const ImagePlaneStrided& dp = dd.plane(0); + const size_t rowBytes = static_cast(sizes[i].w) * channels * sizeof(BT); + for (int32_t y = 0; y < sizes[i].h; ++y) { + std::memcpy(dstHost.data() + static_cast(y) * sizes[i].w * channels, + static_cast(dp.basePtr) + y * dp.rowStride, rowBytes); + } + CompareVectors(dstHost, hostPixels[i]); + } +} + +} // namespace + +int main(int argc, char** argv) { + (void)argc; + (void)argv; + TEST_CASES_BEGIN(); + + // Single-channel, heterogeneous sizes. + TEST_CASE(TestRoundtripCopy({{16, 12}, {32, 24}, {7, 5}, {48, 9}}, FMT_U8)); + TEST_CASE(TestRoundtripCopy({{16, 12}, {32, 24}, {7, 5}, {48, 9}}, FMT_F32)); + + // Multi-channel interleaved. + TEST_CASE(TestRoundtripCopy({{16, 12}, {32, 24}, {7, 5}, {48, 9}}, FMT_RGB8)); + TEST_CASE(TestRoundtripCopy({{16, 12}, {32, 24}, {7, 5}, {48, 9}}, FMT_RGBA8)); + + // Homogeneous batch — the wrapper should still work when all images share the same shape. + TEST_CASE(TestRoundtripCopy({{64, 64}, {64, 64}, {64, 64}}, FMT_RGBA8)); + + // Single image, large. + TEST_CASE(TestRoundtripCopy({{128, 96}}, FMT_RGB8)); + + // CPU path: the same roundtrip driven entirely from host code over a host-resident snapshot. + TEST_CASE(TestHostRoundtripCopy({{16, 12}, {32, 24}, {7, 5}, {48, 9}}, FMT_U8)); + TEST_CASE(TestHostRoundtripCopy({{16, 12}, {32, 24}, {7, 5}, {48, 9}}, FMT_F32)); + TEST_CASE(TestHostRoundtripCopy({{16, 12}, {32, 24}, {7, 5}, {48, 9}}, FMT_RGB8)); + TEST_CASE(TestHostRoundtripCopy({{64, 64}, {64, 64}}, FMT_RGBA8)); + + TEST_CASE(TestAccessors({{16, 12}, {32, 24}, {7, 5}}, FMT_RGB8)); + + // BorderWrapper composed over ImageBatchVarShapeWrapper: constant-fill via guaranteed-OOB read. + TEST_CASE( + TestBorderConstantComposition({{16, 12}, {32, 24}, {7, 5}}, FMT_RGB8, make_uchar3(0xAB, 0xCD, 0xEF))); + TEST_CASE(TestBorderConstantComposition({{16, 12}, {32, 24}, {7, 5}}, FMT_RGBA8, + make_uchar4(0x12, 0x34, 0x56, 0x78))); + + // InterpolationWrapper composed over ImageBatchVarShapeWrapper: integer-coord roundtrip is identity. + TEST_CASE(TestInterpolationNearestComposition({{16, 12}, {32, 24}, {7, 5}}, FMT_U8)); + TEST_CASE(TestInterpolationNearestComposition({{16, 12}, {32, 24}, {7, 5}}, FMT_RGB8)); + TEST_CASE(TestInterpolationNearestComposition({{16, 12}, {32, 24}, {7, 5}}, FMT_RGBA8)); + + TEST_CASES_END(); +} diff --git a/tests/roccv/cpp/src/tests/core/wrappers/test_interpolation_wrapper.cpp b/tests/roccv/cpp/src/tests/core/wrappers/test_interpolation_wrapper.cpp index a4466530..43e07329 100644 --- a/tests/roccv/cpp/src/tests/core/wrappers/test_interpolation_wrapper.cpp +++ b/tests/roccv/cpp/src/tests/core/wrappers/test_interpolation_wrapper.cpp @@ -21,9 +21,9 @@ #include #include -#include "core/detail/vector_utils.hpp" #include +#include "core/detail/vector_utils.hpp" #include "test_helpers.hpp" using namespace roccv; @@ -45,7 +45,7 @@ namespace { * @return T The interpolated pixel. */ template -T GoldenLinear(BorderWrapper input, int64_t sample, float y, float x) { +T GoldenLinear(BorderWrapper> input, int64_t sample, float y, float x) { // Defines the vectorized float type for intermediate calculations. using WorkType = detail::MakeType>; @@ -86,7 +86,7 @@ T GoldenLinear(BorderWrapper input, int64_t sample, float y, floa * @return T The interpolated pixel. */ template -T GoldenNearest(BorderWrapper input, int64_t sample, float y, float x) { +T GoldenNearest(BorderWrapper> input, int64_t sample, float y, float x) { // Nearest neighbor interpolation. Rounds given floating point values to the nearest integer. return input.at(sample, lroundf(y), lroundf(x), 0); } @@ -98,7 +98,7 @@ T GoldenNearest(BorderWrapper input, int64_t sample, float y, flo * @return None. */ void CalBicubicWeights(float dist, float* weight) { - const float A = -0.5f; // Note OpenCV sets alpha to -0.75f + const float A = -0.5f; // Note OpenCV sets alpha to -0.75f weight[0] = ((A * (dist + 1) - 5 * A) * (dist + 1) + 8 * A) * (dist + 1) - 4 * A; weight[1] = ((A + 2) * dist - (A + 3)) * dist * dist + 1; @@ -107,7 +107,8 @@ void CalBicubicWeights(float dist, float* weight) { } /** - * @brief Golden model for Bicubic interpolation. This is the Catmull-Rom cubic interpolation commonly used in CV libraries. + * @brief Golden model for Bicubic interpolation. This is the Catmull-Rom cubic interpolation commonly used in CV + * libraries. * * @tparam T Image datatype. * @tparam BorderType Border type for boundary conditions. @@ -118,7 +119,7 @@ void CalBicubicWeights(float dist, float* weight) { * @return T The interpolated pixel. */ template -T GoldenBicubic(BorderWrapper input, int64_t sample, float y, float x) { +T GoldenBicubic(BorderWrapper> input, int64_t sample, float y, float x) { // Defines the vectorized float type for intermediate calculations. using WorkType = detail::MakeType>; @@ -135,7 +136,8 @@ T GoldenBicubic(BorderWrapper input, int64_t sample, float y, flo WorkType sum = SetAll(0.0f); for (int indexY = -1; indexY <= 2; indexY++) { for (int indexX = -1; indexX <= 2; indexX++) { - sum += detail::RangeCast(input.at(sample, intY + indexY, intX + indexX, 0)) * (weightX[indexX + 1] * weightY[indexY + 1]); + sum += detail::RangeCast(input.at(sample, intY + indexY, intX + indexX, 0)) * + (weightX[indexX + 1] * weightY[indexY + 1]); } } @@ -156,7 +158,7 @@ T GoldenBicubic(BorderWrapper input, int64_t sample, float y, flo * @return T The interpolated pixel. */ template -T GoldenInterpolationAt(BorderWrapper input, int64_t sample, float y, float x, +T GoldenInterpolationAt(BorderWrapper> input, int64_t sample, float y, float x, eInterpolationType interp) { switch (interp) { case eInterpolationType::INTERP_TYPE_NEAREST: @@ -202,9 +204,10 @@ void TestCorrectness(int64_t batchSize, Size2D imageSize, float4 borderValue, fl std::vector> goldenOutput; // Use roccv::InterpolationWrapper to get actual output - InterpolationWrapper actualWrap( - (BorderWrapper(ImageWrapper(input, batchSize, imageSize.w, imageSize.h), borderVal))); - BorderWrapper goldenWrap(ImageWrapper(input, batchSize, imageSize.w, imageSize.h), borderVal); + auto actualWrap = MakeInterpolationWrapper( + MakeBorderWrapper(TensorWrapper(input, batchSize, imageSize.w, imageSize.h), borderVal)); + auto goldenWrap = + MakeBorderWrapper(TensorWrapper(input, batchSize, imageSize.w, imageSize.h), borderVal); for (int b = 0; b < batchSize; b++) { for (float y = 0; y < imageSize.h; y += idxDelta) { @@ -220,7 +223,8 @@ void TestCorrectness(int64_t batchSize, Size2D imageSize, float4 borderValue, fl } } } - if constexpr (std::is_integral_v> && std::is_signed_v> && sizeof(detail::BaseType) == 4) { + if constexpr (std::is_integral_v> && std::is_signed_v> && + sizeof(detail::BaseType) == 4) { CompareVectorsNear(actualOutput, goldenOutput, NEAR_EQUAL_THRESHOLD * 2); } else { CompareVectorsNear(actualOutput, goldenOutput); @@ -228,7 +232,7 @@ void TestCorrectness(int64_t batchSize, Size2D imageSize, float4 borderValue, fl } } // namespace -int main(int argc, char **argv) { +int main(int argc, char** argv) { (void)argc; (void)argv; TEST_CASES_BEGIN(); @@ -322,7 +326,7 @@ int main(int argc, char **argv) { TEST_CASE((TestCorrectness(1, {20, 53}, make_float4(0, 0, 0, 1), 0.1f))); TEST_CASE((TestCorrectness(3, {38, 10}, make_float4(0, 0, 0, 1), 0.1f))); TEST_CASE((TestCorrectness(5, {65, 21}, make_float4(1, 0.5, 0.5, 1), 0.1f))); - // clang-format on + // clang-format on TEST_CASES_END(); } \ No newline at end of file diff --git a/tests/roccv/cpp/src/tests/core/wrappers/test_image_wrapper.cpp b/tests/roccv/cpp/src/tests/core/wrappers/test_tensor_wrapper.cpp similarity index 89% rename from tests/roccv/cpp/src/tests/core/wrappers/test_image_wrapper.cpp rename to tests/roccv/cpp/src/tests/core/wrappers/test_tensor_wrapper.cpp index 70f03b0b..28a1acca 100644 --- a/tests/roccv/cpp/src/tests/core/wrappers/test_image_wrapper.cpp +++ b/tests/roccv/cpp/src/tests/core/wrappers/test_tensor_wrapper.cpp @@ -20,7 +20,7 @@ */ #include -#include +#include #include "test_helpers.hpp" @@ -37,11 +37,11 @@ void TestCorrectness(int numImages, Size2D size) { std::vector ref(numElements); FillVector(ref); - ImageWrapper input(ref, numImages, size.w, size.h); + TensorWrapper input(ref, numImages, size.w, size.h); std::vector actual; // To determine if coordinates are pointing to the proper values in memory, iterate through the reference vector - // element-by-element and iterate through the ImageWrapper coordinate-wise. All values should be the same if + // element-by-element and iterate through the TensorWrapper coordinate-wise. All values should be the same if // everything lines up. for (int b = 0; b < numImages; ++b) { @@ -58,9 +58,9 @@ void TestCorrectness(int numImages, Size2D size) { } template -void TestImageWrapperConstructor(int imageCount, Size2D imageSize, ImageFormat format) { +void TestTensorWrapperConstructor(int imageCount, Size2D imageSize, ImageFormat format) { Tensor input(imageCount, imageSize, format); - ImageWrapper wrapper(input); + TensorWrapper wrapper(input); EXPECT_EQ(wrapper.batches(), imageCount); EXPECT_EQ(wrapper.channels(), format.channels()); @@ -74,7 +74,7 @@ int main(int argc, char** argv) { (void)argv; TEST_CASES_BEGIN(); - TEST_CASE(TestImageWrapperConstructor(2, {54, 67}, FMT_RGB8)); + TEST_CASE(TestTensorWrapperConstructor(2, {54, 67}, FMT_RGB8)); TEST_CASE(TestCorrectness(1, {10, 10})); TEST_CASE(TestCorrectness(2, {43, 9})); diff --git a/tests/roccv/cpp/src/tests/operators/test_op_bilateral_filter.cpp b/tests/roccv/cpp/src/tests/operators/test_op_bilateral_filter.cpp index f208962c..d34a5187 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_bilateral_filter.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_bilateral_filter.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. #include #include #include -#include +#include #include #include "test_helpers.hpp" @@ -51,8 +51,8 @@ namespace { template > void GenerateGoldenBilateral(std::vector& input, std::vector& output, int32_t batchSize, Size2D imageSize, int diameter, float sigmaColor, float sigmaSpace, T borderValue) { - BorderWrapper src(ImageWrapper(input, batchSize, imageSize.w, imageSize.h), borderValue); - ImageWrapper dst(output, batchSize, imageSize.w, imageSize.h); + auto src = MakeBorderWrapper(TensorWrapper(input, batchSize, imageSize.w, imageSize.h), borderValue); + TensorWrapper dst(output, batchSize, imageSize.w, imageSize.h); using namespace roccv::detail; using Worktype = MakeType>; @@ -179,9 +179,9 @@ int main(int argc, char** argv) { TEST_CASE((TestCorrectness(1, 20, 20, FMT_U8, 0, 50.0f, 1.2f, {0.0, 0.0, 0.0, 0.0}, eDeviceType::GPU))); TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGB8, -1, 50.0f, 1.2f, - {0.0, 0.0, 0.0, 0.0}, eDeviceType::GPU))); - TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 0, 500.0f, 1.2f, - {500.0, 500.0, 0.0, 0.0}, eDeviceType::GPU))); + {0.0, 0.0, 0.0, 0.0}, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 0, 500.0f, 1.2f, {500.0, 500.0, 0.0, 0.0}, + eDeviceType::GPU))); TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB8, 4, 50.0f, 3.0f, {0.0, 0.0, 0.0, 0.0}, eDeviceType::GPU))); @@ -288,9 +288,9 @@ int main(int argc, char** argv) { TEST_CASE((TestCorrectness(1, 20, 20, FMT_U8, 0, 50.0f, 1.2f, {0.0, 0.0, 0.0, 0.0}, eDeviceType::CPU))); TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGB8, -1, 50.0f, 1.2f, - {0.0, 0.0, 0.0, 0.0}, eDeviceType::CPU))); - TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 0, 500.0f, 1.2f, - {500.0, 500.0, 0.0, 0.0}, eDeviceType::CPU))); + {0.0, 0.0, 0.0, 0.0}, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 0, 500.0f, 1.2f, {500.0, 500.0, 0.0, 0.0}, + eDeviceType::CPU))); TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB8, 4, 50.0f, 3.0f, {0.0, 0.0, 0.0, 0.0}, eDeviceType::CPU))); diff --git a/tests/roccv/cpp/src/tests/operators/test_op_bnd_box.cpp b/tests/roccv/cpp/src/tests/operators/test_op_bnd_box.cpp index 3b5b5284..3121d7d9 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_bnd_box.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_bnd_box.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include #include -#include +#include #include #include @@ -85,8 +85,8 @@ template > void GenerateGoldenBndBox(std::vector &input, std::vector &output, int32_t batchSize, int32_t width, int32_t height, const BndBoxes &bboxes) { // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); // Working type for internal pixel format, which has 4 channels. using WorkType = detail::MakeType; diff --git a/tests/roccv/cpp/src/tests/operators/test_op_brightness_contrast.cpp b/tests/roccv/cpp/src/tests/operators/test_op_brightness_contrast.cpp index 0532a014..3a9cb8aa 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_brightness_contrast.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_brightness_contrast.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. #include #include -#include +#include #include #include "test_helpers.hpp" @@ -61,8 +61,8 @@ std::vector GoldenBrightnessContrast(std::vector& input, int32_ std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); using work_type = detail::MakeType>; diff --git a/tests/roccv/cpp/src/tests/operators/test_op_center_crop.cpp b/tests/roccv/cpp/src/tests/operators/test_op_center_crop.cpp index f145b1b4..e4e9aa95 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_center_crop.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_center_crop.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include #include -#include +#include #include #include "test_helpers.hpp" @@ -48,8 +48,8 @@ namespace { template > void GenerateGoldenCrop(std::vector& input, std::vector& output, int32_t batchSize, int32_t width, int32_t height, Size2D cropSize) { // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, cropSize.w, cropSize.h); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, cropSize.w, cropSize.h); int topLeftX = (width >> 1) - (cropSize.w >> 1); int topLeftY = (height >> 1) - (cropSize.h >> 1); diff --git a/tests/roccv/cpp/src/tests/operators/test_op_composite.cpp b/tests/roccv/cpp/src/tests/operators/test_op_composite.cpp index 43abf6c1..f332016f 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_composite.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_composite.cpp @@ -21,7 +21,7 @@ #include #include -#include +#include #include #include "test_helpers.hpp" @@ -56,17 +56,17 @@ std::vector> GoldenComposite(std::vector>; - // Wrap input data into ImageWrappers for easy data access - ImageWrapper fgWrap(foreground, batchSize, width, height); - ImageWrapper bgWrap(background, batchSize, width, height); - ImageWrapper maskWrap(mask, batchSize, width, height); + // Wrap input data into TensorWrappers for easy data access + TensorWrapper fgWrap(foreground, batchSize, width, height); + TensorWrapper bgWrap(background, batchSize, width, height); + TensorWrapper maskWrap(mask, batchSize, width, height); // Size of the output depends on the requested number of output channels. If it is 3, then the output images will // have 3 channels. If it is 4, then an additional alpha channel is added to the output. This alpha channel is // always fully on. int numOutElements = batchSize * width * height * detail::NumElements; std::vector> output(numOutElements); - ImageWrapper outWrap(output, batchSize, width, height); + TensorWrapper outWrap(output, batchSize, width, height); for (int b = 0; b < batchSize; b++) { for (int y = 0; y < height; y++) { diff --git a/tests/roccv/cpp/src/tests/operators/test_op_convert_to.cpp b/tests/roccv/cpp/src/tests/operators/test_op_convert_to.cpp index 6f14bdbd..cb7fb3d9 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_convert_to.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_convert_to.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. #include #include -#include +#include #include #include "test_helpers.hpp" @@ -55,8 +55,8 @@ std::vector GoldenConvertTo(std::vector& input, int32_t batchSi std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); using AB_DT = decltype(float() * BT_SRC() * BT_DEST()); using work_type = detail::MakeType>; diff --git a/tests/roccv/cpp/src/tests/operators/test_op_copy_make_border.cpp b/tests/roccv/cpp/src/tests/operators/test_op_copy_make_border.cpp index 4320f04e..34a0191a 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_copy_make_border.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_copy_make_border.cpp @@ -57,10 +57,11 @@ std::vector GoldenCopyMakeBorder(std::vector input, int batchSize, Size2 // Wrap the input images in a BorderWrapper to handle out of bounds image behavior. The BorderWrapper has already // been tested in another test so it can be used reliably. - BorderWrapper inputWrap(ImageWrapper(input, batchSize, inputSize.w, inputSize.h), borderVal); + auto inputWrap = + MakeBorderWrapper(TensorWrapper(input, batchSize, inputSize.w, inputSize.h), borderVal); std::vector output(batchSize * outputSize.h * outputSize.w * channels); - ImageWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); + TensorWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); for (int b = 0; b < batchSize; b++) { for (int y = 0; y < outputSize.h; y++) { diff --git a/tests/roccv/cpp/src/tests/operators/test_op_custom_crop.cpp b/tests/roccv/cpp/src/tests/operators/test_op_custom_crop.cpp index 03cbf9b4..a161cea7 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_custom_crop.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_custom_crop.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. #include #include #include -#include +#include #include #include @@ -50,8 +50,8 @@ template > void GenerateGoldenCrop(std::vector& input, std::vector& output, int32_t batchSize, int32_t width, int32_t height, Box_t cropRect) { // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, cropRect.width, cropRect.height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, cropRect.width, cropRect.height); for (int b = 0; b < batchSize; b++) { for (int y = 0; y < cropRect.height; y++) { diff --git a/tests/roccv/cpp/src/tests/operators/test_op_cvt_color.cpp b/tests/roccv/cpp/src/tests/operators/test_op_cvt_color.cpp index 52c3735d..3c06c397 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_cvt_color.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_cvt_color.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. #include #include -#include +#include #include #include "test_helpers.hpp" @@ -34,9 +34,9 @@ namespace { template > std::vector GoldenReorder(std::vector& input, int samples, int width, int height) { - ImageWrapper inputWrap(input, samples, width, height); + TensorWrapper inputWrap(input, samples, width, height); std::vector output(samples * width * height * detail::NumElements); - ImageWrapper outputWrap(output, samples, width, height); + TensorWrapper outputWrap(output, samples, width, height); for (int b = 0; b < samples; b++) { for (int y = 0; y < height; y++) { @@ -51,9 +51,9 @@ std::vector GoldenReorder(std::vector& input, int samples, int width, in template > std::vector GoldenYUVToRGB(std::vector& input, int samples, int width, int height, float delta) { - ImageWrapper inputWrap(input, samples, width, height); + TensorWrapper inputWrap(input, samples, width, height); std::vector output(samples * width * height * detail::NumElements); - ImageWrapper outputWrap(output, samples, width, height); + TensorWrapper outputWrap(output, samples, width, height); for (int b = 0; b < samples; b++) { for (int y = 0; y < height; y++) { @@ -77,9 +77,9 @@ std::vector GoldenYUVToRGB(std::vector& input, int samples, int width, i template > std::vector GoldenRGBToYUV(std::vector& input, int samples, int width, int height, float delta) { - ImageWrapper inputWrap(input, samples, width, height); + TensorWrapper inputWrap(input, samples, width, height); std::vector output(samples * width * height * detail::NumElements); - ImageWrapper outputWrap(output, samples, width, height); + TensorWrapper outputWrap(output, samples, width, height); for (int b = 0; b < samples; b++) { for (int y = 0; y < height; y++) { @@ -104,11 +104,11 @@ std::vector GoldenRGBToYUV(std::vector& input, int samples, int width, i template > std::vector GoldenRGBToGrayscale(std::vector& input, int samples, int width, int height) { - ImageWrapper inputWrap(input, samples, width, height); + TensorWrapper inputWrap(input, samples, width, height); std::vector output(samples * width * height); // Output must always be uchar1 for grayscale - ImageWrapper outputWrap(output, samples, width, height); + TensorWrapper outputWrap(output, samples, width, height); for (int b = 0; b < samples; b++) { for (int y = 0; y < height; y++) { diff --git a/tests/roccv/cpp/src/tests/operators/test_op_flip.cpp b/tests/roccv/cpp/src/tests/operators/test_op_flip.cpp index 36bfbcc4..1a83da71 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_flip.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_flip.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. #include #include #include -#include +#include #include #include "test_helpers.hpp" @@ -52,8 +52,8 @@ std::vector GoldenFlip(std::vector& input, int32_t batchSize, int32_t wi std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { for (int y = 0; y < height; ++y) { diff --git a/tests/roccv/cpp/src/tests/operators/test_op_gamma_contrast.cpp b/tests/roccv/cpp/src/tests/operators/test_op_gamma_contrast.cpp index d24df72e..71b1014e 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_gamma_contrast.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_gamma_contrast.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" #include "core/detail/math/vectorized_type_math.hpp" -#include +#include #include #include #include "operator_types.h" @@ -56,8 +56,8 @@ std::vector GoldenGammaContrast(std::vector& input, int32_t batchSize, i std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); using work_type = detail::MakeType>; for (int b = 0; b < batchSize; ++b) { diff --git a/tests/roccv/cpp/src/tests/operators/test_op_histogram.cpp b/tests/roccv/cpp/src/tests/operators/test_op_histogram.cpp index 05a5009c..513df448 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_histogram.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_histogram.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ #include -#include +#include #include #include #include @@ -57,7 +57,7 @@ std::vector GoldenHistogram(std::vector& input, int32_t batchSize, in std::vector local_histogram(256); // Wrap the input vector for simplified data access - ImageWrapper src(input, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { std::fill(local_histogram.begin(), local_histogram.end(), 0); @@ -94,8 +94,8 @@ std::vector GoldenHistogramMask(std::vector& input, std::vector local_histogram(256); // Wrap input/mask vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper maskWrap(mask, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper maskWrap(mask, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { std::fill(local_histogram.begin(), local_histogram.end(), 0); diff --git a/tests/roccv/cpp/src/tests/operators/test_op_normalize.cpp b/tests/roccv/cpp/src/tests/operators/test_op_normalize.cpp index fd850ddb..3b077b36 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_normalize.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_normalize.cpp @@ -22,7 +22,7 @@ THE SOFTWARE. #include #include -#include +#include #include #include "test_helpers.hpp" diff --git a/tests/roccv/cpp/src/tests/operators/test_op_remap.cpp b/tests/roccv/cpp/src/tests/operators/test_op_remap.cpp index 634344a4..18847056 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_remap.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_remap.cpp @@ -21,12 +21,13 @@ THE SOFTWARE. */ #include -#include +#include #include #include #include -#include "core/detail/internal_structs.hpp" + #include "core/detail/casting.hpp" +#include "core/detail/internal_structs.hpp" #include "core/detail/math/vectorized_type_math.hpp" #include "core/detail/type_traits.hpp" #include "operator_types.h" @@ -39,11 +40,11 @@ using namespace roccv::detail; // Keep all non-entrypoint functions in an anonymous namespace to prevent redefinition errors across translation units. namespace { -RemapParams GetRemapParams(const int2 &srcSize, const int2 &dstSize, const int2 &mapSize, bool alignCorners, eRemapType mapValueType) -{ +RemapParams GetRemapParams(const int2& srcSize, const int2& dstSize, const int2& mapSize, bool alignCorners, + eRemapType mapValueType) { RemapParams params; - switch(mapValueType) { + switch (mapValueType) { case REMAP_ABSOLUTE: params.srcScale = make_float2(0.f, 0.f); params.mapScale = StaticCast(mapSize) / StaticCast(dstSize); @@ -54,7 +55,7 @@ RemapParams GetRemapParams(const int2 &srcSize, const int2 &dstSize, const int2 case REMAP_ABSOLUTE_NORMALIZED: params.srcScale = make_float2(0.f, 0.f); params.mapScale = StaticCast(mapSize) / StaticCast(dstSize); - params.valScale = (StaticCast(srcSize) - (alignCorners ? 1.f : 0.f)) / 2.f; + params.valScale = (StaticCast(srcSize) - (alignCorners ? 1.f : 0.f)) / 2.f; params.srcOffset = params.valScale - (alignCorners ? 0.f : .5f); params.dstOffset = 0.f; break; @@ -87,24 +88,24 @@ RemapParams GetRemapParams(const int2 &srcSize, const int2 &dstSize, const int2 */ template > -std::vector GoldenRemap(std::vector& input, int32_t batchSize, int32_t mapBatchSize, int32_t inWidth, int32_t inHeight, int32_t outWidth, - int32_t outHeight, int32_t mapWidth, int32_t mapHeight, std::vector& mapData, eRemapType mapType, bool alignCorners, float4 borderValue) { - +std::vector GoldenRemap(std::vector& input, int32_t batchSize, int32_t mapBatchSize, int32_t inWidth, + int32_t inHeight, int32_t outWidth, int32_t outHeight, int32_t mapWidth, int32_t mapHeight, + std::vector& mapData, eRemapType mapType, bool alignCorners, float4 borderValue) { int channels = detail::NumElements; int outputSize = batchSize * outWidth * outHeight * channels; std::vector output(outputSize); // Create interpolation wrapper for input vector - InterpolationWrapper src((BorderWrapper( - ImageWrapper(input, batchSize, inWidth, inHeight), detail::SaturateCast(borderValue)))); + auto src = MakeInterpolationWrapper(MakeBorderWrapper( + TensorWrapper(input, batchSize, inWidth, inHeight), detail::SaturateCast(borderValue))); // Wrap the output vector for simplified data access - ImageWrapper dst(output, batchSize, outWidth, outHeight); + TensorWrapper dst(output, batchSize, outWidth, outHeight); // Create an interpolation wrapper for the map tensor - // InterpolationWrapper wrappedMapTensor(map, make_float2(0, 0)); - InterpolationWrapper map((BorderWrapper( - ImageWrapper(mapData.data(), mapBatchSize, mapWidth, mapHeight), detail::SaturateCast(borderValue)))); + auto map = MakeInterpolationWrapper( + MakeBorderWrapper(TensorWrapper(mapData.data(), mapBatchSize, mapWidth, mapHeight), + detail::SaturateCast(borderValue))); int2 srcSize = make_int2(src.width(), src.height()); int2 dstSize = make_int2(dst.width(), dst.height()); @@ -119,13 +120,12 @@ std::vector GoldenRemap(std::vector& input, int32_t batchSize, int32_t m for (int b = 0; b < dst.batches(); b++) { for (int y = 0; y < dst.height(); y++) { for (int x = 0; x < dst.width(); x++) { - dstCoord.x = static_cast(x); dstCoord.y = static_cast(y); - + mapCoord.x = (dstCoord.x + params.dstOffset) * params.mapScale.x; mapCoord.y = (dstCoord.y + params.dstOffset) * params.mapScale.y; - + float2 mapValue = map.at((mapBatchSize == 1 ? 0 : b), mapCoord.y, mapCoord.x, 0); srcCoord.x = dstCoord.x * params.srcScale.x + mapValue.x * params.valScale.x + params.srcOffset.x; @@ -162,7 +162,8 @@ std::vector GoldenRemap(std::vector& input, int32_t batchSize, int32_t m */ template > -void TestCorrectness(int batchSize, int mapBatchSize, int inWidth, int inHeight, int outWidth, int outHeight, int mapWidth, int mapHeight, ImageFormat format, float4 borderValue, eRemapType mapType, +void TestCorrectness(int batchSize, int mapBatchSize, int inWidth, int inHeight, int outWidth, int outHeight, + int mapWidth, int mapHeight, ImageFormat format, float4 borderValue, eRemapType mapType, bool alignCorners, eDeviceType device) { // Create input and output tensor based on test parameters Tensor input(batchSize, {inWidth, inHeight}, format, device); @@ -174,7 +175,7 @@ void TestCorrectness(int batchSize, int mapBatchSize, int inWidth, int inHeight, // Copy generated input data into input tensor CopyVectorIntoTensor(input, inputData); - + int mapSize = mapBatchSize * mapWidth * mapHeight; std::vector mapData(mapSize); @@ -188,11 +189,10 @@ void TestCorrectness(int batchSize, int mapBatchSize, int inWidth, int inHeight, } } } - } - else if (mapType == REMAP_ABSOLUTE_NORMALIZED) { + } else if (mapType == REMAP_ABSOLUTE_NORMALIZED) { for (int b = 0; b < mapBatchSize; b++) { - for (int y = 0; y < mapHeight; y++){ - for (int x = 0; x < mapWidth; x++){ + for (int y = 0; y < mapHeight; y++) { + for (int x = 0; x < mapWidth; x++) { float normX = ((2.0f * static_cast(x)) / static_cast(mapWidth - 1)) - 1.0f; float normY = ((2.0f * static_cast(y)) / static_cast(mapHeight - 1)) - 1.0f; @@ -204,11 +204,10 @@ void TestCorrectness(int batchSize, int mapBatchSize, int inWidth, int inHeight, } } } - } - else if (mapType == REMAP_RELATIVE_NORMALIZED) { + } else if (mapType == REMAP_RELATIVE_NORMALIZED) { for (int b = 0; b < mapBatchSize; b++) { - for (int y = 0; y < mapHeight; y++){ - for (int x = 0; x < mapWidth; x++){ + for (int y = 0; y < mapHeight; y++) { + for (int x = 0; x < mapWidth; x++) { // Generate normalized coordinates in [-1, 1] range float normX = ((2.0f * static_cast(x)) / static_cast(mapWidth - 1)) - 1.0f; float normY = ((2.0f * static_cast(y)) / static_cast(mapHeight - 1)) - 1.0f; @@ -235,7 +234,8 @@ void TestCorrectness(int batchSize, int mapBatchSize, int inWidth, int inHeight, hipStream_t stream; HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream)); Remap op; - op(stream, input, output, mapTensor, InterpType, MapInterpType, mapType, alignCorners, BorderType, borderValue, device); + op(stream, input, output, mapTensor, InterpType, MapInterpType, mapType, alignCorners, BorderType, borderValue, + device); HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream)); HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream)); @@ -243,9 +243,9 @@ void TestCorrectness(int batchSize, int mapBatchSize, int inWidth, int inHeight, std::vector result(output.shape().size()); CopyTensorIntoVector(result, output); - std::vector ref = GoldenRemap(inputData, batchSize, mapBatchSize, inWidth, - inHeight, outWidth, outHeight, - mapWidth, mapHeight, mapData, mapType, alignCorners, borderValue); + std::vector ref = GoldenRemap( + inputData, batchSize, mapBatchSize, inWidth, inHeight, outWidth, outHeight, mapWidth, mapHeight, mapData, + mapType, alignCorners, borderValue); // Compare data in actual output versus the generated golden reference image CompareVectors(result, ref); @@ -258,144 +258,186 @@ int main(int argc, char** argv) { TEST_CASES_BEGIN(); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, false, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, false, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, false, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, false, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, false, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, false, eDeviceType::GPU))); - + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + false, eDeviceType::GPU))); + TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, true, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, true, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, true, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, true, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, true, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, true, eDeviceType::GPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::GPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::GPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, false, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, false, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, false, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, false, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, false, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, false, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, true, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, true, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, true, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, true, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGB8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, true, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE_NORMALIZED, + true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, true, eDeviceType::CPU))); + 1, 1, 480, 360, 480, 360, 480, 360, FMT_RGBA8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_RELATIVE_NORMALIZED, + true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, false, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, false, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::CPU))); + eInterpolationType::INTERP_TYPE_NEAREST>(2, 1, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::CPU))); TEST_CASE((TestCorrectness( - 2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, make_float4(0.0f, 0.0f, 0.0f, 1.0f), REMAP_ABSOLUTE, true, eDeviceType::CPU))); - - + eInterpolationType::INTERP_TYPE_NEAREST>(2, 2, 480, 360, 480, 360, 480, 360, FMT_U8, + make_float4(0.0f, 0.0f, 0.0f, 1.0f), + REMAP_ABSOLUTE, true, eDeviceType::CPU))); TEST_CASES_END(); } \ No newline at end of file diff --git a/tests/roccv/cpp/src/tests/operators/test_op_resize.cpp b/tests/roccv/cpp/src/tests/operators/test_op_resize.cpp index d7c385d0..61ed04d7 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_resize.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_resize.cpp @@ -20,10 +20,18 @@ 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 #include +#include #include "test_helpers.hpp" @@ -45,18 +53,17 @@ namespace { * @return A vector containing the data for the images resized to outputSize. */ template > -std::vector GoldenResize(std::vector> &input, int batchSize, Size2D inputSize, +std::vector GoldenResize(std::vector>& input, int batchSize, Size2D inputSize, Size2D outputSize) { size_t numOutputElements = batchSize * outputSize.w * outputSize.h * detail::NumElements; std::vector> output(numOutputElements); - ImageWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); + TensorWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); // Use the replicate (or clamping) border mode by default to handle out of bounds conditions with certain // interpolation modes. - InterpolationWrapper inputWrap( - BorderWrapper( - ImageWrapper(input, batchSize, inputSize.w, inputSize.h), T{})); + auto inputWrap = MakeInterpolationWrapper(MakeBorderWrapper( + TensorWrapper(input, batchSize, inputSize.w, inputSize.h), T{})); // Determine the scaling factor required to map from the output coordinates to the corresponding input coordinates // on both the x and y axes. @@ -126,9 +133,96 @@ void TestCorrectness(int batchSize, Size2D inputSize, Size2D outputSize, ImageFo CompareVectorsNear(actualOutput, goldenOutput); } +/** + * @brief Copies a packed host pixel buffer into a single-plane Image, honoring the plane's row stride and the + * image's residency. + */ +template > +void CopyVectorIntoImage(Image& img, const std::vector& src, Size2D size, eDeviceType device) { + const int channels = detail::NumElements; + const size_t rowBytes = static_cast(size.w) * channels * sizeof(BT); + + switch (device) { + case eDeviceType::GPU: { + auto data = img.exportData(); + const ImagePlaneStrided& p = data.plane(0); + HIP_VALIDATE_NO_ERRORS( + hipMemcpy2D(p.basePtr, p.rowStride, src.data(), rowBytes, rowBytes, size.h, hipMemcpyHostToDevice)); + break; + } + case eDeviceType::CPU: { + auto data = img.exportData(); + const ImagePlaneStrided& p = data.plane(0); + for (int32_t y = 0; y < size.h; ++y) { + std::memcpy(static_cast(p.basePtr) + y * p.rowStride, + src.data() + static_cast(y) * size.w * channels, rowBytes); + } + break; + } + } +} + +/** + * @brief Compares roccv::Resize over a variable-shape image batch against a per-image golden model. + * + * Each input image has its own size; all are resized into a single constant-sized output tensor (NHWC). The golden + * result reuses GoldenResize per image (batch of 1, that image's input size) and concatenates, matching the NHWC + * output ordering. + * + * @tparam T The image's pixel datatype. + * @tparam InterpType The interpolation type to use during resizing. + * @param inputSizes Per-image input sizes (defines the batch size). + * @param outputSize The uniform output size for every image. + * @param format The image format to use (must match with the given type T). + * @param device The device to run this correctness test on. + */ +template > +void TestCorrectnessVarShape(const std::vector& inputSizes, Size2D outputSize, ImageFormat format, + eDeviceType device) { + const int channels = detail::NumElements; + const int32_t numImages = static_cast(inputSizes.size()); + + // Build a variable-shape batch from per-image random host data. + ImageBatchVarShape batch(numImages, device); + std::vector images; + images.reserve(numImages); + std::vector> hostInputs(numImages); + + for (int32_t i = 0; i < numImages; ++i) { + images.emplace_back(inputSizes[i], format, device); + hostInputs[i].resize(static_cast(inputSizes[i].w) * inputSizes[i].h * channels); + FillVector(hostInputs[i], static_cast(0x500 + i)); + CopyVectorIntoImage(images[i], hostInputs[i], inputSizes[i], device); + batch.pushBack(images[i]); + } + + // Uniform, constant-sized output tensor. + Tensor outputTensor(numImages, outputSize, format, device); + + hipStream_t stream; + HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream)); + Resize op; + op(stream, batch, outputTensor, InterpType, device); + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream)); + HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream)); + + std::vector actualOutput(outputTensor.shape().size()); + CopyTensorIntoVector(actualOutput, outputTensor); + + // Per-image golden, concatenated in batch order to match the NHWC output layout. + std::vector goldenOutput; + goldenOutput.reserve(actualOutput.size()); + for (int32_t i = 0; i < numImages; ++i) { + std::vector g = GoldenResize(hostInputs[i], 1, inputSizes[i], outputSize); + goldenOutput.insert(goldenOutput.end(), g.begin(), g.end()); + } + + CompareVectorsNear(actualOutput, goldenOutput); +} + } // namespace -int main(int argc, char **argv) { +int main(int argc, char** argv) { (void)argc; (void)argv; TEST_CASES_BEGIN(); @@ -157,6 +251,16 @@ int main(int argc, char **argv) { TEST_CASE((TestCorrectness(3, {100, 50}, {100, 50}, FMT_RGBf32, eDeviceType::GPU))); TEST_CASE((TestCorrectness(5, {100, 50}, {50, 25}, FMT_RGBAf32, eDeviceType::GPU))); + // U8 - Cubic interpolation + TEST_CASE((TestCorrectness(1, {100, 50}, {200, 50}, FMT_U8, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(3, {100, 50}, {100, 50}, FMT_RGB8, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(5, {100, 50}, {50, 25}, FMT_RGBA8, eDeviceType::GPU))); + + // F32 - Cubic interpolation + TEST_CASE((TestCorrectness(1, {100, 50}, {200, 50}, FMT_F32, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(3, {100, 50}, {100, 50}, FMT_RGBf32, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(5, {100, 50}, {50, 25}, FMT_RGBAf32, eDeviceType::GPU))); + // CPU Tests // U8 - Linear interpolation @@ -178,6 +282,37 @@ int main(int argc, char **argv) { TEST_CASE((TestCorrectness(1, {100, 50}, {200, 50}, FMT_F32, eDeviceType::CPU))); TEST_CASE((TestCorrectness(3, {100, 50}, {100, 50}, FMT_RGBf32, eDeviceType::CPU))); TEST_CASE((TestCorrectness(5, {100, 50}, {50, 25}, FMT_RGBAf32, eDeviceType::CPU))); + + // U8 - Cubic interpolation + TEST_CASE((TestCorrectness(1, {100, 50}, {200, 50}, FMT_U8, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(3, {100, 50}, {100, 50}, FMT_RGB8, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(5, {100, 50}, {50, 25}, FMT_RGBA8, eDeviceType::CPU))); + + // F32 - Cubic interpolation + TEST_CASE((TestCorrectness(1, {100, 50}, {200, 50}, FMT_F32, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(3, {100, 50}, {100, 50}, FMT_RGBf32, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(5, {100, 50}, {50, 25}, FMT_RGBAf32, eDeviceType::CPU))); + + // Variable-shape batch -> constant-sized tensor. Heterogeneous input sizes (upscale, downscale, and odd + // dimensions) into a single uniform output. + + // GPU + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_U8, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGB8, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGBA8, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGB8, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_F32, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGBf32, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGBAf32, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGB8, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGBf32, eDeviceType::GPU))); + + // CPU + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_U8, eDeviceType::CPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGB8, eDeviceType::CPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGBA8, eDeviceType::CPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGBf32, eDeviceType::CPU))); + TEST_CASE((TestCorrectnessVarShape({{100, 50}, {37, 91}, {200, 13}}, {64, 64}, FMT_RGB8, eDeviceType::CPU))); // clang-format on TEST_CASES_END(); diff --git a/tests/roccv/cpp/src/tests/operators/test_op_rotate.cpp b/tests/roccv/cpp/src/tests/operators/test_op_rotate.cpp index 56deeabb..fb9fa55e 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_rotate.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_rotate.cpp @@ -67,10 +67,9 @@ std::vector> GoldenRotate(std::vector>& T borderVal = detail::SaturateCast(make_float4(0.0f, 0.0f, 0.0f, 0.0f)); - ImageWrapper outputWrapper(output, batchSize, imageSize.w, imageSize.h); - InterpolationWrapper inputWrapper( - BorderWrapper(ImageWrapper(input, batchSize, imageSize.w, imageSize.h), - borderVal)); + TensorWrapper outputWrapper(output, batchSize, imageSize.w, imageSize.h); + auto inputWrapper = MakeInterpolationWrapper(MakeBorderWrapper( + TensorWrapper(input, batchSize, imageSize.w, imageSize.h), borderVal)); /** * Affine warp for a combined rotation and translate looks like the following when in its inverse representation: diff --git a/tests/roccv/cpp/src/tests/operators/test_op_thresholding.cpp b/tests/roccv/cpp/src/tests/operators/test_op_thresholding.cpp index 09ef3262..279959e0 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_thresholding.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_thresholding.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. #include "core/detail/casting.hpp" #include "core/detail/type_traits.hpp" #include "core/detail/math/vectorized_type_math.hpp" -#include +#include #include #include #include "operator_types.h" @@ -59,8 +59,8 @@ std::vector GoldenBinaryThreshold(std::vector& input, int32_t batchSize, std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { double th = thresh[b]; @@ -88,8 +88,8 @@ std::vector GoldenBinaryInvThreshold(std::vector& input, int32_t batchSi std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { double th = thresh[b]; @@ -117,8 +117,8 @@ std::vector GoldenTruncThreshold(std::vector& input, int32_t batchSize, std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { double th = thresh[b]; @@ -145,8 +145,8 @@ std::vector GoldenToZeroThreshold(std::vector& input, int32_t batchSize, std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { double th = thresh[b]; @@ -173,8 +173,8 @@ std::vector GoldenToZeroInvThreshold(std::vector& input, int32_t batchSi std::vector output(input.size()); // Wrap input/output vectors for simplified data access - ImageWrapper src(input, batchSize, width, height); - ImageWrapper dst(output, batchSize, width, height); + TensorWrapper src(input, batchSize, width, height); + TensorWrapper dst(output, batchSize, width, height); for (int b = 0; b < batchSize; ++b) { double th = thresh[b]; diff --git a/tests/roccv/cpp/src/tests/operators/test_op_warp_affine.cpp b/tests/roccv/cpp/src/tests/operators/test_op_warp_affine.cpp index 93c91ae9..5af99149 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_warp_affine.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_warp_affine.cpp @@ -55,12 +55,12 @@ std::vector> GoldenWarpAffine(std::vector& mat, bool isInverted, int batchSize, Size2D inputSize, Size2D outputSize, float4 borderValue) { // Create interpolation wrapper for input vector - InterpolationWrapper inputWrap((BorderWrapper( - ImageWrapper(input, batchSize, inputSize.w, inputSize.h), detail::SaturateCast(borderValue)))); + auto inputWrap = MakeInterpolationWrapper(MakeBorderWrapper( + TensorWrapper(input, batchSize, inputSize.w, inputSize.h), detail::SaturateCast(borderValue))); - // Create ImageWrapper for output vector. We also need to create said output vector. + // Create TensorWrapper for output vector. We also need to create said output vector. std::vector> output(batchSize * outputSize.w * outputSize.h * detail::NumElements); - ImageWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); + TensorWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); // Prepare the transformation matrix. An affine transform is effectively a 3x3 perspective transform with its last // row set to [0, 0, 1]. diff --git a/tests/roccv/cpp/src/tests/operators/test_op_warp_perspective.cpp b/tests/roccv/cpp/src/tests/operators/test_op_warp_perspective.cpp index 1461365c..2c7559cb 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_warp_perspective.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_warp_perspective.cpp @@ -52,12 +52,12 @@ std::vector> GoldenWarpPerspective(std::vector& mat, bool isInverted, int batchSize, Size2D inputSize, Size2D outputSize, float4 borderValue) { // Create interpolation wrapper for input vector - InterpolationWrapper inputWrap((BorderWrapper( - ImageWrapper(input, batchSize, inputSize.w, inputSize.h), detail::SaturateCast(borderValue)))); + auto inputWrap = MakeInterpolationWrapper(MakeBorderWrapper( + TensorWrapper(input, batchSize, inputSize.w, inputSize.h), detail::SaturateCast(borderValue))); - // Create ImageWrapper for output vector. We also need to create said output vector. + // Create TensorWrapper for output vector. We also need to create said output vector. std::vector> output(batchSize * outputSize.w * outputSize.h * detail::NumElements); - ImageWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); + TensorWrapper outputWrap(output, batchSize, outputSize.w, outputSize.h); // If given matrix is not the inverted representation of the transformation, we have to invert it first (since we // transform from output -> input).