Skip to content

Architecture: Kernel Wrappers

Jump to bottom
Zachary Vincze edited this page Jun 1, 2026 · 1 revision

Kernel Wrapper Chain

This page documents the kernel wrapper chain in include/core/wrappers/: the non-owning accessor types that bridge containers and kernels, how they compose, and how operators build them.

What the chain is

A kernel wrapper is a small, trivially copyable, non-owning accessor over container memory. Operators build a wrapper on the host, pass it by value into a kernel, and the kernel reads/writes pixels through it. Wrappers carry pointers and stride math — never ownership — so they free nothing.

Wrappers compose. A base accessor over a container is decorated by a border-handling wrapper, which is decorated by an interpolation wrapper. Each layer adds one responsibility and forwards the rest:

base accessor → BorderWrapper → InterpolationWrapper → consumed by kernel

The composition works because every layer satisfies the same informal wrapper concept:

Member Meaning
ValueType Pixel type the wrapper reads/writes (e.g. uchar3, float4)
at(n, h, w, c) Access channel c of pixel (h, w) in image n
width(n) / height(n) Per-image dimensions; n defaults to 0
batches() Number of images
channels() Channels per pixel

Because the border and interpolation decorators are templated on the wrapper they wrap (not on a fixed base type), the same border/interpolation math serves both uniform tensors and variable-shape image batches with no duplication.

Base accessors

Each container has exactly one base accessor. Both expose the wrapper concept; they differ in how they resolve at, width, and height.

TensorWrapper<T>

include/core/wrappers/tensor_wrapper.hpp — bounds-unsafe accessor over a Tensor with NHWC, NCHW, HWC, or CHW layout.

  • Template: T — pixel type. BaseType = detail::BaseType<T> is the scalar element.
  • Construction: from a Tensor (reads layout, normalizes to {n, h, w, c} shape/stride; HWC/CHW get batch shape 1, batch stride 0); also from a std::vector<BaseType>& or a void* with explicit batchSize/width/height.
  • at(n, h, w, c) computes a byte offset from per-dimension strides and reinterprets to T&. No bounds checking.
  • width(n) / height(n) ignore n and return the single shared dimension. The unused n parameter exists only so the signature matches ImageBatchVarShapeWrapper, letting both satisfy the same concept (include/core/wrappers/tensor_wrapper.hpp:149).
  • Other layouts throw NOT_IMPLEMENTED at construction.

Used for both input and output. Output tensors use it directly (no border/interpolation decoration).

ImageBatchVarShapeWrapper<T>

include/core/wrappers/image_batch_var_shape_wrapper.hpp — accessor over an ImageBatchVarShape's exported descriptor table.

  • Template: T — single pixel type (uchar1, uchar3, uchar4, float1, float4, …). Channel count derives from detail::NumElements<T>.
  • Construction: from an ImageBatchVarShapeDataStrided snapshot (input.exportData(stream)). Copies out the descriptor-table pointer and image count.
  • at(n, h, w, c) indexes descriptor n, then applies single-plane interleaved (NHWC-style) stride math: basePtr + h*rowStride + w*sizeof(T) + c*sizeof(BaseType).
  • width(n) / height(n) read per-image dimensions from descriptor n. This is the difference that makes variable shapes work.
  • Residency follows the snapshot. A GPU (...Hip) snapshot yields a device pointer (use in device kernels); a CPU (...Host) snapshot yields a host pointer (use in host code). Every at/width/height dereferences it, so the caller must run the wrapper on the matching side.
  • Single-plane interleaved only; ImageBatchVarShape rejects multi-plane images at pushBack.

GenericTensorWrapper<T>

include/core/wrappers/generic_tensor_wrapper.hpp — rank-agnostic accessor for tensors that are not images.

  • Template: T — element type.
  • Construction: from a Tensor (copies shape/strides/rank), or from raw void* + shape + strides + rank.
  • at(idx...) takes a variadic coordinate list and dots it against strides.
  • Does not expose width/height/batches/channels, so it does not satisfy the image wrapper concept and does not compose with BorderWrapper/InterpolationWrapper. It is for side-channel data such as histogram bins or threshold tables. See op_histogram.cpp and thresholding_host.hpp.

Border decorator

include/core/wrappers/border_wrapper.hpp — adds out-of-bounds handling on top of a base accessor.

template <eBorderType BorderType, typename W>
class BorderWrapper;
  • Templates: BorderType — the border policy (compile-time); W — the wrapped base accessor type. ValueType is recovered from W::ValueType; the border mode is exposed as the kBorderType constant and W as WrapperType.
  • Construction: BorderWrapper(W image_wrapper, ValueType border_value). border_value is the fallback used only by BORDER_TYPE_CONSTANT.
  • at(n, h, w, c) returns the in-bounds value when (w, h) lies inside width(n) × height(n); otherwise it remaps the coordinate per the border policy (or returns the constant) and reads through W.

In-bounds reads take an early-return branch to skip the remap math for the common case (include/core/wrappers/border_wrapper.hpp:87).

Border types

From include/operator_types.h (eBorderType):

Value Behavior
BORDER_TYPE_CONSTANT Returns border_value for out-of-bounds coordinates
BORDER_TYPE_REPLICATE Clamps the coordinate to the last valid row/column
BORDER_TYPE_REFLECT Reflects across the edge, including the boundary pixel
BORDER_TYPE_REFLECT101 Reflects across the edge, excluding the boundary pixel
BORDER_TYPE_WRAP Wraps the coordinate to the opposite edge

Each policy is a if constexpr branch selected at compile time, so only the chosen one is instantiated.

Factory

template <eBorderType B, typename W>
auto MakeBorderWrapper(W wrap, typename W::ValueType borderValue);

Deduces W from the argument so callers spell only the border-mode policy.

Interpolation decorator

include/core/wrappers/interpolation_wrapper.hpp — adds sub-pixel sampling on top of a BorderWrapper. Read-only; do not use for output tensors.

template <eInterpolationType I, typename BW>
class InterpolationWrapper;
  • Templates: I — interpolation policy (compile-time); BW — the BorderWrapper type. ValueType comes from BW::ValueType; BW is exposed as BorderType and I as kInterpolationType. Templating directly on the BorderWrapper means the border mode and base wrapper are spelled once.
  • Construction: InterpolationWrapper(BW borderWrapper).
  • at(n, h, w, c) takes float h/w and samples. Reads go through the inner BorderWrapper::at, so out-of-bounds taps are border-handled automatically.

Interpolation types

From include/operator_types.h (eInterpolationType):

Value Sampling
INTERP_TYPE_NEAREST Rounds to the nearest pixel (lroundf)
INTERP_TYPE_LINEAR Bilinear over the 2×2 neighborhood
INTERP_TYPE_CUBIC Catmull-Rom bicubic over the 4×4 neighborhood

Linear and cubic promote each tap to a float work type (detail::MakeType<float, NumElements<ValueType>>) via detail::RangeCast, accumulate, then range-cast back to ValueType. Cubic weights come from CalBicubicWeights.

Factory

template <eInterpolationType I, typename BW>
auto MakeInterpolationWrapper(BW borderWrap);

Deduces BW (and its border mode + base wrapper) from the argument; callers spell only the interpolation policy.

Usage: the Resize operator

src/op_resize.cpp builds the chain once and reuses it for both the uniform-tensor and variable-shape input paths. The base wrapper is built by the caller; everything downstream is identical (src/op_resize.cpp:42).

template <typename T, eInterpolationType I, typename SrcBaseWrapper>
void dispatch_resize_launch(hipStream_t stream, SrcBaseWrapper srcBase,
                            const Tensor& output, eDeviceType device) {
    TensorWrapper<T> outputWrapper(output);

    // Resize clamps at the image border (REPLICATE).
    auto inputWrapper = MakeInterpolationWrapper<I>(
        MakeBorderWrapper<eBorderType::BORDER_TYPE_REPLICATE>(srcBase, T{}));

    switch (device) {
        case eDeviceType::GPU: {
            dim3 block(64, 16);
            dim3 grid(/* over output width/height */, outputWrapper.batches());
            Kernels::Device::resize<<<grid, block, 0, stream>>>(inputWrapper, outputWrapper);
            break;
        }
        case eDeviceType::CPU:
            Kernels::Host::resize(inputWrapper, outputWrapper);
            break;
    }
}

The two entry points differ only in the base wrapper they construct (src/op_resize.cpp:65):

// Uniform tensor input
dispatch_resize_launch<T, I>(stream, TensorWrapper<T>(input), output, device);

// Variable-shape batch input
auto data = input.exportData(stream);
dispatch_resize_launch<T, I>(stream, ImageBatchVarShapeWrapper<T>(data), output, device);

The kernel is wrapper-agnostic — it names no concrete wrapper type, only the concept (include/kernels/device/resize_device.hpp):

template <typename SrcWrapper, typename DstWrapper>
__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<float>(output.width());
    float scaleY = input.height(batch) / static_cast<float>(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);
}

input.width(batch) resolves to the shared width for TensorWrapper and to image batch's own width for ImageBatchVarShapeWrapper — the per-image scale falls out for free. The host kernel (include/kernels/host/resize_host.hpp) is the same logic under an OpenMP loop.

Reference

Type Header Role
TensorWrapper<T> include/core/wrappers/tensor_wrapper.hpp Base accessor for NHWC/NCHW/HWC/CHW tensors
ImageBatchVarShapeWrapper<T> include/core/wrappers/image_batch_var_shape_wrapper.hpp Base accessor for variable-shape image batches
GenericTensorWrapper<T> include/core/wrappers/generic_tensor_wrapper.hpp Rank-agnostic accessor for non-image data
BorderWrapper<BorderType, W> include/core/wrappers/border_wrapper.hpp Border-handling decorator + MakeBorderWrapper
InterpolationWrapper<I, BW> include/core/wrappers/interpolation_wrapper.hpp Interpolation decorator + MakeInterpolationWrapper
eBorderType, eInterpolationType include/operator_types.h Border and interpolation policy enums

Clone this wiki locally