Skip to content

Contributing: Writing Tests

Jump to bottom
Zachary Vincze edited this page May 28, 2026 · 2 revisions

Writing Tests

Every new operator needs a C++ test and a Python test. This page walks through what to write, using Flip as the running example. For running tests, see Contributing: Running Tests.

What to Test

Three categories, in order of importance:

  1. Correctness — the operator produces the expected output for each supported (dtype, channels, layout, device) combination.
  2. Negative cases — the operator throws the expected exception when given invalid input.
  3. Edge sizes — small, large, non-aligned dimensions; batches of 1 and many.

You don't need exhaustive coverage of every combination. Hit the dispatcher branches and the obvious edges.

C++ Tests

Each operator gets its own executable at tests/roccv/cpp/src/tests/operators/test_op_<name>.cpp. The file has its own main() and uses the custom test framework (no gtest).

File Structure

#include <op_flip.hpp>
#include "test_helpers.hpp"

using namespace roccv;
using namespace roccv::tests;

namespace {  // Anonymous — prevents symbol collisions across test binaries

template <typename T>
std::vector<T> GoldenFlip(...) { /* CPU reference */ }

template <typename T>
void TestCorrectness(int batchSize, int width, int height, ...,
                     ImageFormat format, eDeviceType device) { /* ... */ }

void TestNegativeFlip() { /* throw-checks */ }

}  // namespace

int main(int argc, char** argv) {
    (void)argc; (void)argv;
    TEST_CASES_BEGIN();

    TEST_CASE(TestNegativeFlip());
    TEST_CASE(TestCorrectness<uchar3>(1, 480, 360, 0, FMT_RGB8, eDeviceType::GPU));
    TEST_CASE(TestCorrectness<uchar3>(1, 480, 360, 0, FMT_RGB8, eDeviceType::CPU));
    // ... one TEST_CASE per (dtype, channels, device) combination

    TEST_CASES_END();
}

Always wrap helpers in an anonymous namespace — test binaries get linked separately but share symbol space at the CMake level.

Writing the Golden Model

A golden model is a simple, obviously-correct CPU implementation that the kernel is compared against. Write it for clarity, not speed. Use TensorWrapper<T> so indexing matches the kernel's view of memory.

template <typename T, typename BT = detail::BaseType<T>>
std::vector<BT> GoldenFlip(std::vector<BT>& input, int batchSize, int w, int h, int flipCode) {
    std::vector<BT> output(input.size());
    TensorWrapper<T> src(input, batchSize, w, h);
    TensorWrapper<T> dst(output, batchSize, w, h);

    for (int b = 0; b < batchSize; ++b)
        for (int y = 0; y < h; ++y)
            for (int x = 0; x < w; ++x) {
                int srcX = flipCode > 0 ? (w - 1 - x) : x;
                int srcY = flipCode <= 0 ? (h - 1 - y) : y;
                if (flipCode < 0) { srcX = w - 1 - x; srcY = h - 1 - y; }
                dst.at(b, y, x, 0) = detail::SaturateCast<T>(src.at(b, srcY, srcX, 0));
            }
    return output;
}

Rules of thumb:

  • Triple-nested loops over (b, y, x) are fine — clarity beats cleverness.
  • Use detail::SaturateCast<T> when the math might overflow a fixed-point type.
  • Don't share code between the golden model and the kernel. The whole point is independent implementations.

The Correctness Harness

A reusable template that runs the operator and compares against the golden model:

template <typename T, typename BT = detail::BaseType<T>>
void TestCorrectness(int batchSize, int w, int h, int flipCode,
                     ImageFormat format, eDeviceType device) {
    Tensor input(batchSize, {w, h}, format, device);
    Tensor output(batchSize, {w, h}, format, device);

    std::vector<BT> inputData(input.shape().size());
    FillVector(inputData);                       // random fill
    CopyVectorIntoTensor(input, inputData);

    std::vector<BT> ref = GoldenFlip<T>(inputData, batchSize, w, h, flipCode);

    hipStream_t stream;
    HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream));
    Flip op;
    op(stream, input, output, flipCode, device);
    HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream));
    HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream));

    std::vector<BT> result(output.shape().size());
    CopyTensorIntoVector(result, output);
    CompareVectors(result, ref);
}

Test helpers like FillVector, CopyVectorIntoTensor, CompareVectors, and HIP_VALIDATE_NO_ERRORS live in tests/roccv/cpp/include/test_helpers.hpp.

Exact vs. Near Comparison

Two comparison helpers — pick based on whether your operator does floating-point math.

Helper When to use
CompareVectors(result, ref) Bit-exact match required. Integer-only ops, pure copy/permutation kernels (Flip, CopyMakeBorder), or anything where you can guarantee identical math on both paths.
CompareVectorsNear(result, ref, delta) Floating-point math, or integer math that goes through an FP intermediate (resize, rotate, convert with scaling, color-space transforms). 
CompareVectorsNear(result, ref);              // default delta = 1e-6 (normalized)
CompareVectorsNear(result, ref, 1e-4);        // looser — for cubic interp, large rotations, etc.

GPU vs. CPU Float Precision

GPU and CPU paths will not produce bit-identical floating-point output. Sources of divergence:

  • FMA on GPU, separate mul+add on CPU — different rounding behavior.
  • Transcendentals (sinf, cosf, expf) use different polynomial approximations.
  • Reduction order in any kernel doing per-pixel accumulation differs from the CPU loop order.
  • -ffast-math-style optimizations in the device compiler.

This is expected and unavoidable. Use CompareVectorsNear for any operator touching floats.

Picking delta — it's specified in the normalized [0, 1] range; the helper scales it to the target dtype via RangeCast<T> (and clamps to at least 1 for integers, so an F32 delta of 1e-6 becomes a tolerance of 1 for U8).

Operator class Suggested delta
Pure copy/reshape (no math) Use CompareVectors (exact)
Linear arithmetic, small integer accumulators 1e-6 (default)
Linear interpolation (resize, warp) 1e-5 to 1e-4
Cubic interpolation, color conversion 1e-4 to 1e-3
Trig-heavy ops (rotate by arbitrary angle) 1e-3

Tighten, don't loosen. Start with 1e-6 and raise the delta only when you've confirmed the failure is precision-related (not a real bug). A loose delta hides regressions.

If GPU and CPU paths share floating-point structure but still drift past 1e-3, that's a smell — investigate before bumping the threshold further.

Negative Tests

For each CHECK_* macro in the operator, write one test case that triggers it. Use EXPECT_EXCEPTION to assert the right eStatusType is thrown.

void TestNegativeFlip() {
    TensorShape shape(TensorLayout(TENSOR_LAYOUT_NHWC), {1, 1, 1, 1});
    Tensor gpuTensor(shape, DataType(DATA_TYPE_U8), eDeviceType::GPU);
    Tensor cpuTensor(shape, DataType(DATA_TYPE_U8), eDeviceType::CPU);
    Flip op;

    // Output tensor on wrong device
    EXPECT_EXCEPTION(op(nullptr, gpuTensor, cpuTensor, 0, eDeviceType::GPU),
                     eStatusType::INVALID_OPERATION);

    // Unsupported layout
    Tensor badLayout(TensorShape(TensorLayout(TENSOR_LAYOUT_NC), {1, 1}),
                     DataType(DATA_TYPE_U8), eDeviceType::GPU);
    EXPECT_EXCEPTION(op(nullptr, badLayout, gpuTensor, 0, eDeviceType::GPU),
                     eStatusType::INVALID_COMBINATION);

    // Unsupported dtype
    Tensor badDtype(shape, DataType(DATA_TYPE_U32), eDeviceType::GPU);
    EXPECT_EXCEPTION(op(nullptr, badDtype, gpuTensor, 0, eDeviceType::GPU),
                     eStatusType::NOT_IMPLEMENTED);

    // Shape mismatch
    Tensor wrongShape(TensorShape(gpuTensor.layout(), {2, 2, 2, 2}),
                      DataType(DATA_TYPE_U8), eDeviceType::GPU);
    EXPECT_EXCEPTION(op(nullptr, wrongShape, gpuTensor, 0, eDeviceType::GPU),
                     eStatusType::INVALID_COMBINATION);
}

Picking Test Cases

For each operator, at minimum test:

  • Every supported dtype (U8, S32, F32, ...)
  • Every supported channel count (1, 3, 4)
  • Both devices (CPU and GPU) — they must both produce correct output when compared to the golden output
  • A few odd dimensions (e.g. 134 × 360, 23 × 106) to catch alignment bugs
  • A batch > 1 to catch batch-stride bugs

Use the TEST_CASE(...) macro one line per combination. Don't loop — explicit lines make failures grep-able.

Python Tests

Python tests live at tests/roccv/python/test_op_<name>.py and use pytest's parametrization to fan out combinations.

import pytest
import rocpycv
from test_helpers import generate_tensor, compare_tensors


@pytest.mark.parametrize("device", [rocpycv.eDeviceType.CPU, rocpycv.eDeviceType.GPU])
@pytest.mark.parametrize("dtype",  [rocpycv.eDataType.U8, rocpycv.eDataType.S32, rocpycv.eDataType.F32])
@pytest.mark.parametrize("channels", [1, 3, 4])
@pytest.mark.parametrize("flip_code", [-1, 0, 1])
@pytest.mark.parametrize("samples,width,height", [
    (1, 65, 30),
    (3, 45, 20),
    (20, 104, 234),
])
def test_op_flip(samples, width, height, channels, dtype, flip_code, device):
    input_tensor = generate_tensor(samples, width, height, channels, dtype, device)
    stream = rocpycv.Stream()

    # Test the in-place ("_into") variant against the allocating variant
    out_golden = rocpycv.Tensor([samples, height, width, channels],
                                rocpycv.eTensorLayout.NHWC, dtype, device)
    rocpycv.flip_into(out_golden, input_tensor, flip_code, stream, device)

    out = rocpycv.flip(input_tensor, flip_code, stream, device)
    stream.synchronize()

    compare_tensors(out, out_golden)

What the Python Test Covers

The Python test isn't a duplicate of the C++ one — it verifies the binding works:

  • Both variants exist: allocating (flip) and in-place (flip_into).
  • All parametrized dtype/channel/device combinations dispatch correctly through pybind11.
  • The inferred output tensor shape for the "non-into" version of the operator matches an expected output shape.

Heavy correctness logic belongs in the C++ test. The Python test ensures the binding plumbing is intact.

Helpers (generate_tensor, compare_tensors) live in tests/roccv/python/test_helpers.py.

Checklist Before Pushing

  • Every supported (dtype, channels) combination has a TEST_CASE line
  • Both eDeviceType::CPU and eDeviceType::GPU covered
  • One negative test per CHECK_* macro in the operator
  • Golden model is independent (no shared helpers with the kernel)
  • Python test parametrizes dtype, channels, device, and exercises both API variants
  • ctest -R test_op_<name> and pytest test_op_<name>.py both pass

Next Steps

Clone this wiki locally