Skip to content

Contributing: Writing a New Operator

Jump to bottom
Zachary Vincze edited this page May 27, 2026 · 5 revisions

Contributing a New Operator

This page walks through the structure of a rocCV operator: what files to create, how to document the API, validate inputs, dispatch on type/layout, and launch kernels.

Shortcut: the new-operator skill (Claude Code) scaffolds all eight files and wires registration from a single spec. This page documents the structure so you can write or review one by hand.

Files You'll Touch

File Role
include/op_<name>.hpp Public operator class (inherits IOperator)
src/op_<name>.cpp Validation + dispatcher + kernel launch
include/kernels/device/<name>_device.hpp __global__ GPU kernel(s)
include/kernels/host/<name>_host.hpp CPU kernel(s)
tests/roccv/cpp/src/tests/operators/test_op_<name>.cpp C++ test
tests/roccv/python/test_op_<name>.py Python test
python/include/operators/Op<Name>.h + python/src/operators/Op<Name>.cpp pybind11 binding
include/roccv_operators.hpp Add #include "op_<name>.hpp"
python/src/main.cpp Register PyOp<Name>::Export(m)

CMake uses GLOB_RECURSE — no CMakeLists.txt edits needed.

1. Public Header & Documentation

The header declares the class and carries the user-facing contract. Document supported layouts, dtypes, channel counts, and the input↔output relationship. This block is what users read.

/**
 * @brief Flips an image batch about the horizontal, vertical, or both axes.
 *
 * Input:
 *       Supported TensorLayout(s): [NHWC, HWC]
 *                        Channels: [1, 3, 4]
 *       Supported DataType(s):     [U8, S32, F32]
 *
 * Output:
 *       (same as input)
 *
 * Input/Output dependency:
 *       Property      |  Input == Output
 *      ---------------|-----------------
 *       TensorLayout  | Yes
 *       DataType      | Yes
 *       Shape         | Yes
 *
 * @param[in]  stream    HIP stream.
 * @param[in]  input     Input image batch.
 * @param[out] output    Output image batch.
 * @param[in]  flipCode  0 = X-axis, >0 = Y-axis, <0 = both.
 * @param[in]  device    GPU (default) or CPU.
 */
void operator()(hipStream_t stream, const Tensor& input, const Tensor& output,
                int32_t flipCode, eDeviceType device = eDeviceType::GPU) const;

Documentation rules:

  • List supported layouts/dtypes/channels explicitly — don't say "any".
  • Include the I/O dependency table even when everything matches; it sets reader expectations.
  • Annotate every parameter with @param[in] or @param[out].
  • See include/op_flip.hpp for a reference template.

2. Input Validation

The first thing operator() does. Use the macros from include/common/validation_helpers.hpp — they throw Exception with the right eStatusType and message.

Macro Checks
CHECK_TENSOR_DEVICE(t, device) Tensor lives on the device we're dispatching to
CHECK_TENSOR_LAYOUT(t, ...) Layout is in the allowed list
CHECK_TENSOR_DATATYPES(t, ...) Dtype is in the allowed list
CHECK_TENSOR_CHANNELS(t, ...) Channel count is in the allowed list
CHECK_TENSOR_COMPARISON(expr) Arbitrary boolean (use for input↔output equality)

Typical block:

CHECK_TENSOR_DEVICE(input, device);
CHECK_TENSOR_DATATYPES(input, DATA_TYPE_U8, DATA_TYPE_S32, DATA_TYPE_F32);
CHECK_TENSOR_CHANNELS(input, 1, 3, 4);
CHECK_TENSOR_LAYOUT(input, TENSOR_LAYOUT_HWC, TENSOR_LAYOUT_NHWC);

CHECK_TENSOR_DEVICE(output, device);
CHECK_TENSOR_COMPARISON(input.layout() == output.layout());
CHECK_TENSOR_COMPARISON(input.dtype()  == output.dtype());
CHECK_TENSOR_COMPARISON(input.shape()  == output.shape());

Validate every assumption your kernel makes. The kernel runs untyped at the bit level — a missed check becomes silent corruption.

3. Dispatcher Tables

Operators are templated on dtype, channel count, and any compile-time parameters. The dispatcher chooses the right instantiation at runtime via lookup tables keyed on eDataType and channel count.

Layered pattern (recommended):

  1. Outer table → picks the C++ scalar type (uchar1/uchar3/uchar4, etc.) based on dtype + channels.
  2. Inner dispatcher → resolves any operator-specific compile-time params (flip axis, border type, ...).
  3. Innermost dispatcher → builds wrappers and launches the kernel.

Example outer table (one row per dtype, indexed by channels - 1):

static const std::unordered_map<eDataType, std::array<DispatchFn, 4>> funcs = {
    {DATA_TYPE_U8,  {dispatch<uchar1>, 0, dispatch<uchar3>, dispatch<uchar4>}},
    {DATA_TYPE_S32, {dispatch<int1>,   0, dispatch<int3>,   dispatch<int4>}},
    {DATA_TYPE_F32, {dispatch<float1>, 0, dispatch<float3>, dispatch<float4>}},
};

auto func = funcs.at(input.dtype().etype())[input.shape(input.layout().channels_index()) - 1];
if (func == 0) throw Exception("Not mapped.", eStatusType::INVALID_OPERATION);
func(stream, input, output, ..., device);

The 0 entries mark unsupported channel counts (e.g. 2-channel) — they cause a clean throw rather than UB.

See src/op_flip.cpp for the full pattern.

4. Kernel Launching

The deepest dispatcher builds typed wrappers and dispatches to GPU or CPU.

template <typename T, eAxis FlipType>
void dispatch_kernel(hipStream_t stream, const Tensor& input, const Tensor& output,
                     eDeviceType device) {
    TensorWrapper<T> inputWrapper(input);
    TensorWrapper<T> outputWrapper(output);

    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::flip<FlipType><<<grid, block, 0, stream>>>(
                inputWrapper, outputWrapper);
            break;
        }
        case eDeviceType::CPU: {
            Kernels::Host::flip<FlipType>(inputWrapper, outputWrapper);
            break;
        }
    }
}

Grid sizing: standard pattern is (width × height × batch) — one thread per output pixel, with blockIdx.z covering the batch dimension.

Wrappers (in include/core/wrappers/):

  • TensorWrapper<T> — direct bounds-unsafe accessor.
  • BorderWrapper<T, BT> — adds border handling (replicate, constant, reflect, wrap).
  • InterpolationWrapper<T, BT, IT> — adds interpolation (nearest, linear, cubic) on top of a border wrapper.

Compose them when your kernel reads out-of-bounds or sub-pixel coordinates.

Kernel files are header-only because __global__ templates must be instantiated from the .cpp translation unit that launches them.

5. Tests

Every operator needs both a C++ test and a Python test. The standard pattern is golden-model comparison: run the same op on GPU and CPU, compare outputs. See Contributing: Running Tests for the framework details.

Registration

Two one-line edits:

// include/roccv_operators.hpp
#include "op_<name>.hpp"
// python/src/main.cpp  (inside PYBIND11_MODULE)
PyOp<Name>::Export(m);

Checklist

  • Public header with full Doxygen contract (layouts, dtypes, channels, I/O dependency)
  • All input assumptions validated via CHECK_* macros
  • Dispatcher covers every supported (dtype, channel) combination
  • GPU and CPU paths both implemented
  • C++ test using EXPECT_TEST_STATUS + golden-model comparison
  • Python test mirrors C++ coverage
  • pybind11 binding registered in main.cpp
  • #include added to roccv_operators.hpp

Next Steps

Clone this wiki locally