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/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).