Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
69 changes: 69 additions & 0 deletions include/common/validation_helpers.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,7 @@ THE SOFTWARE.
#include <algorithm>
#include <vector>

#include "core/image_format.hpp"
#include "core/tensor.hpp"

/**
Expand Down Expand Up @@ -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<eDataType> 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<int32_t> 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);
76 changes: 46 additions & 30 deletions include/core/wrappers/border_wrapper.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
Expand Down Expand Up @@ -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<T>, ImageBatchVarShapeWrapper<T>) 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 <typename T, eBorderType BorderType>
template <eBorderType BorderType, typename W>
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<T> 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);
}

Expand All @@ -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);
Expand All @@ -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.)
Expand Down Expand Up @@ -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.
Expand All @@ -213,7 +215,21 @@ class BorderWrapper {
__device__ __host__ inline int64_t channels() const { return m_desc.channels(); }

private:
ImageWrapper<T> m_desc;
T m_border_value;
W m_desc;
ValueType m_border_value;
};
} // namespace roccv

/**
* @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 <eBorderType B, typename W>
auto MakeBorderWrapper(W wrap, typename W::ValueType borderValue) {
return BorderWrapper<B, W>(wrap, borderValue);
}

} // namespace roccv
119 changes: 119 additions & 0 deletions include/core/wrappers/image_batch_var_shape_wrapper.hpp
Original file line number Diff line number Diff line change
@@ -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 <hip/hip_runtime.h>

#include <cassert>
#include <cstdint>

#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<T>
* (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<T>.
*
* 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 <typename T>
class ImageBatchVarShapeWrapper {
public:
using ValueType = T;
using BaseType = detail::BaseType<T>;

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<T>; }

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<T>::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<unsigned char*>(p.basePtr) + h * p.rowStride + w * sizeof(T) + c * sizeof(BaseType);
return reinterpret_cast<T*>(addr);
}

const ImageBufferStrided* m_imageList = nullptr;
int32_t m_numImages = 0;
};

} // namespace roccv
Loading