rocCV Operator Expansion

Goal

Close the operator-parity gap between AMD's rocCV (21 ops on develop) and competitor by landing 10 high-leverage operators end-to-end (HIP kernel + CPU fallback + C++ API + rocpycv binding + tests + benchmarks + docs) over 12 weeks.

1. Summary

rocCV is AMD's GPU-accelerated computer vision operator library. It is the direct counterpart CV-CUDA and is the API surface that PyTorch / vision developers use for on-GPU pre- and post-processing (Resize, Normalize, ColorJitter, etc.). Today it has 21 operators on the develop branch (the public v0.2.0 docs still list only 16). CV-CUDA has 45+ operators. Every missing operator forces a CPU fallback, breaks the zero-copy DLPack story, and makes AMD GPUs uncompetitive for standard PyTorch CV preprocessing pipelines.

This project lands 10 operators chosen from the top of the "actually-used in real PyTorch pipelines" list (BrightnessContrast, GaussianBlur, MedianBlur, AverageBlur, AdaptiveThreshold, MorphologyEx, Pad, RandomResizedCrop, HQResize, HistogramEq + CLAHE). Each is shipped with:

HIP device kernel
Host (CPU) fallback kernel
Public C++ class API matching the existing roccv::Flip-style pattern
rocpycv Python binding with DLPack zero-copy
Unit tests (numerical parity vs OpenCV / torchvision golden data)
Microbenchmark in benchmarks/
Doxygen + Sphinx documentation page

Why this is high-impact. When this project lands, torchvision.transforms.v2 Compose pipelines, Albumentations DSLs, and CV-CUDA-style preprocessing scripts can be ported to AMD with a one-line import change. This is the precondition for follow-on work — for example, registering rocCV as a true torchvision.transforms.v2 backend.

2. Background

2.1 What rocCV is

From the rocCV README:

rocCV is an efficient GPU-accelerated library for image pre- and post-processing, powered by AMD's HIP platform.

Layered architecture:

   PyTorch / NumPy user code
              │  DLPack zero-copy
              ▼
        rocpycv (Python)
              │  pybind11
              ▼
         libroccv (C++)
              │
     ┌────────┴────────┐
     ▼                 ▼
 HIP device         host CPU
 kernels            kernels

2.2 What's already in `rocCV develop` (21 operators)

Verified by listing the upstream tree on develop:

Image / tensor	Geometry	Color	Stats / structural
BilateralFilter	Resize	CvtColor	Histogram
Composite	Rotate	AdvCvtColor	NonMaximumSuppression
CopyMakeBorder	WarpAffine	GammaContrast	Threshold
CustomCrop	WarpPerspective	Normalize
CenterCrop	Remap
BndBox	Flip
ConvertTo (dtype/scale)	Reformat (layout)

Note: The published docs at rocm.docs.amd.com/projects/rocCV/en/latest/reference/rocCV-supported-operators.html are pinned at v0.2.0 and only show 16 operators (Flip, CenterCrop, ConvertTo, Reformat, AdvCvtColor are not listed there but exist on develop). One of this project's stretch deliverables is to refresh that page.

2.3 The gap vs CV-CUDA

CV-CUDA exposes 45+ operators. The gaps that hurt real PyTorch users today (counted by "appears in torchvision.transforms.v2.Compose([...]) and Albumentations pipelines"):

Brightness / Contrast — no fused op (you have GammaContrast which is different math)
Gaussian blur — missing
Average / box blur — missing
Median blur — missing
Adaptive threshold — missing (only global threshold)
Morphology (Erode / Dilate / Open / Close) — missing entirely
Pad / PadStack — only CopyMakeBorder (single-tensor) exists
RandomResizedCrop — missing (torchvision's most-used training op)
High-quality resize (HQResize) — Resize exists but limited interpolation
HistogramEq + CLAHE — Histogram exists but no equalization

This list is what this project ships.

3. Problem statement

A PyTorch user writing this standard preprocessing pipeline today on ROCm:

import torchvision.transforms.v2 as T
preprocess = T.Compose([
    T.Resize(256), T.RandomResizedCrop(224),
    T.RandomHorizontalFlip(),
    T.ColorJitter(brightness=0.2, contrast=0.2),
    T.GaussianBlur(kernel_size=3),
    T.ToDtype(torch.float32, scale=True),
    T.Normalize(mean=IMAGENET_MEAN, std=IMAGENET_STD),
])

…falls back to CPU torchvision for RandomResizedCrop, ColorJitter (which uses brightness/contrast under the hood), and GaussianBlur. The end-to-end data-loader throughput is dominated by these CPU ops. The same code on an NVIDIA box dispatches everything to GPU via CV-CUDA.

Quantitative target. By end of projetct, the above pipeline should run end-to-end on GPU on gfx942 (MI300X) with zero CPU fallbacks for steps 1–6, and ≥ 2× pipeline throughput vs the CPU-torchvision baseline at batch size 256, image size 224×224.

4. Goals and non-goals

4.1 Goals (must-have)

G1. Land 10 new operators in ROCm/rocCV develop via separate PRs (one PR per operator).
G2. Each operator has parity (within tolerance) with its OpenCV / torchvision reference on a generated golden-data test matrix.
G3. Each operator has a HIP device kernel and a CPU host kernel (matching the existing pattern).
G4. Each operator has a rocpycv Python binding with DLPack zero-copy.
G5. Each operator has a microbenchmark in benchmarks/ and an entry in analyze_results.py output.
G6. Each operator has a Doxygen header + Sphinx page; the supported-operators table in docs/ is updated.
G7. Final week: end-to-end ImageNet-pipeline benchmark on PyTorch showing ≥ 2× throughput uplift.

4.2 Non-goals (explicitly out of scope)

❌ Building a torchvision.transforms.v2 dispatch backend (that's the natural follow-on project).
❌ Adding new datatypes beyond {U8, F32} unless the operator's reference requires it (e.g. Histogram needs U8 only).
❌ Modifying roccv::Tensor or rocpycv core infrastructure unless strictly necessary for a chosen operator.
❌ Multi-GPU support, distributed pipelines, or rocAL integration.
❌ Audio or video operators (those are RPP/rocAL territory).

5. Definition of done

For each individual operator, all of the following must be true before the PR can be merged:

For the whole project, additionally:

End-to-end ImageNet preprocess benchmark exists and shows ≥ 2× throughput uplift over CPU torchvision.
Public docs page rocCV-supported-operators lists all 31 operators (21 existing + 10 new).
One blog-style write-up (≤ 1500 words) suitable for posting on the ROCm blog.

6. Target operator set

Each row gives: priority, complexity (1–5), what CV-CUDA / OpenCV / torchvision call it, the reference math, and a one-line "why."

#	Operator	Pri	Cmplx	Ref impl	Why
1	BrightnessContrast	P0	1	`cv::convertScaleAbs`, `torchvision F.adjust_brightness`+`F.adjust_contrast`	Single most-used color aug; fused multiply-add per pixel.
2	GaussianBlur	P0	2	`cv::GaussianBlur`	Used in CLIP/DINO augmentation, segmentation. Separable conv.
3	AverageBlur	P0	1	`cv::blur`	Separable, simpler than Gaussian; pre-req for many compose pipelines.
4	MedianBlur	P1	3	`cv::medianBlur`	Common; non-separable; sort-network kernel on 3×3 / 5×5.
5	AdaptiveThreshold	P1	2	`cv::adaptiveThreshold`	Document & OCR preprocessing; builds on existing Threshold.
6	MorphologyEx	P1	3	`cv::morphologyEx` (erode/dilate/open/close/gradient)	5 ops in one; binary classical-CV staple.
7	Pad / PadStack	P0	1	`torch.nn.functional.pad`, `torchvision F.pad`	Batch padding to common shape; trivial extension of CopyMakeBorder.
8	RandomResizedCrop	P0	2	`torchvision F.resized_crop`	THE ImageNet training transform; combined sample-and-resize.
9	HQResize	P2	3	CV-CUDA's HQ Resize (Lanczos, anti-aliased)	Higher-quality resize for inference; outpaces basic bilinear.
10	HistogramEq + CLAHE	P1	3	`cv::equalizeHist`, `cv::createCLAHE`	Tone mapping; building on existing Histogram operator.

Priority tiers:

P0 (must ship): BrightnessContrast, GaussianBlur, AverageBlur, Pad, RandomResizedCrop — these cover the ImageNet training pipeline and are referenced in the success criterion.
P1 (should ship): MedianBlur, AdaptiveThreshold, MorphologyEx, HistogramEq + CLAHE.
P2 (nice to have): HQResize.

7. Codebase tour

The upstream layout on develop:

rocCV/
├── include/
│   ├── i_operator.hpp                # Base interface every op implements
│   ├── operator_types.h              # Enums: eDeviceType, eInterpolationType, etc.
│   ├── roccv_operators.hpp           # Umbrella header — add new op_*.hpp include here
│   ├── op_<snake>.hpp                # Public op header (one per op)
│   ├── core/                         # Tensor, TensorShape, TensorLayout, etc.
│   ├── common/                       # Shared utilities
│   └── kernels/
│       ├── device/
│       │   └── <snake>_device.hpp    # HIP __global__ kernel + launcher
│       ├── host/
│       │   └── <snake>_host.hpp      # CPU reference / fallback kernel
│       ├── common/                   # Shared kernel helpers
│       └── kernel_helpers.hpp        # Block / grid / index helpers
├── src/
│   ├── CMakeLists.txt                # Register new op_*.cpp here
│   ├── op_<snake>.cpp                # Dispatch CPU vs GPU; validate args
│   └── core/                         # Tensor implementation
├── python/
│   ├── CMakeLists.txt
│   ├── include/                      # Python binding helpers
│   └── src/
│       ├── main.cpp                  # Module entry; register op submodule
│       ├── py_tensor.cpp             # Tensor bindings
│       ├── py_enums.cpp
│       ├── operators/
│       │   └── py_op_<snake>.cpp     # pybind11 bindings (one per op)
│       └── rocpycv.pyi               # Type stubs — add Foo class
├── tests/
│   ├── cpp/                          # C++ unit tests (ctest)
│   └── pybind/                       # Python tests (pytest)
├── benchmarks/
│   ├── config.json                   # Bench config; add new op
│   ├── analyze_results.py
│   ├── roccvbench/                   # Bench harness
│   └── src/                          # Per-op bench source
├── samples/                          # End-to-end demos
└── docs/
    └── reference/                    # Sphinx; add operators/<snake>.rst

Read these first (in this order):

include/op_flip.hpp and src/op_flip.cpp — simplest representative op
include/kernels/device/flip_device.hpp — HIP kernel pattern
include/i_operator.hpp — base interface
include/core/tensor.hpp — tensor abstraction (Requirements, exportData, dtype, layout)
python/src/operators/py_op_flip.cpp (assuming the convention; otherwise look for Flip in main.cpp) — binding pattern
tests/cpp/test_op_flip.cpp — testing pattern
benchmarks/src/bench_op_flip.cpp — benchmark pattern

8. The operator implementation template

Use Flip as the canonical 1-input → 1-output operator template. BrightnessContrast (operator #1 below) follows exactly this shape and is the recommended first ticket.

8.1 Public header — `include/op_brightness_contrast.hpp`

#pragma once

#include <operator_types.h>
#include <i_operator.hpp>
#include "core/tensor.hpp"

namespace roccv {
/**
 * @brief Adjusts brightness and contrast of every image in a batch.
 *
 * For each pixel: out = saturate(alpha * in + beta)
 *   - alpha = contrast factor (typically [0, 2])
 *   - beta  = brightness offset (typically [-128, 128] for U8)
 *
 * Limitations:
 *   Input/Output layouts:  NHWC, HWC
 *   Channels:              1, 3, 4
 *   Datatypes:             U8, F32
 *   Input == Output in shape, layout, datatype, channels.
 *
 * @param[in]  stream HIP stream.
 * @param[in]  input  Input tensor (image batch).
 * @param[out] output Output tensor.
 * @param[in]  alpha  Contrast factor per image (scalar or [N]-shaped tensor).
 * @param[in]  beta   Brightness offset per image (scalar or [N]-shaped tensor).
 * @param[in]  device GPU (default) or CPU.
 */
class BrightnessContrast final : public IOperator {
 public:
    BrightnessContrast() = default;
    ~BrightnessContrast() = default;

    void operator()(hipStream_t stream,
                    const Tensor& input,
                    const Tensor& output,
                    const Tensor& alpha,
                    const Tensor& beta,
                    eDeviceType device = eDeviceType::GPU) const;
};
}  // namespace roccv

8.2 Dispatch — `src/op_brightness_contrast.cpp`

#include "op_brightness_contrast.hpp"
#include "kernels/device/brightness_contrast_device.hpp"
#include "kernels/host/brightness_contrast_host.hpp"

namespace roccv {
void BrightnessContrast::operator()(hipStream_t stream,
                                    const Tensor& input,
                                    const Tensor& output,
                                    const Tensor& alpha,
                                    const Tensor& beta,
                                    eDeviceType device) const {
    // 1. Validate: same shape, layout, dtype for input/output.
    // 2. Validate: dtype in {U8, F32}; channels in {1,3,4}.
    // 3. Validate: alpha, beta shape == [N] or scalar.
    // 4. Dispatch:
    if (device == eDeviceType::GPU) {
        device_kernels::launch_brightness_contrast(stream, input, output, alpha, beta);
    } else {
        host_kernels::run_brightness_contrast(input, output, alpha, beta);
    }
}
}  // namespace roccv

8.3 HIP device kernel — `include/kernels/device/brightness_contrast_device.hpp`

#pragma once
#include <hip/hip_runtime.h>
#include "core/tensor.hpp"
#include "kernel_helpers.hpp"

namespace roccv::device_kernels {

template <typename T>
__global__ void k_brightness_contrast(const T* __restrict__ in,
                                      T* __restrict__ out,
                                      const float* alpha,
                                      const float* beta,
                                      int N, int H, int W, int C) {
    int n = blockIdx.z;
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x >= W || y >= H || n >= N) return;

    const float a = alpha[n];
    const float b = beta[n];
    int base = ((n * H + y) * W + x) * C;
    #pragma unroll
    for (int c = 0; c < C; ++c) {
        float v = static_cast<float>(in[base + c]) * a + b;
        out[base + c] = saturate_cast<T>(v);  // see kernels/common/saturate.hpp
    }
}

void launch_brightness_contrast(hipStream_t stream,
                                const Tensor& input, const Tensor& output,
                                const Tensor& alpha, const Tensor& beta);
}  // namespace roccv::device_kernels

8.4 Host fallback — `include/kernels/host/brightness_contrast_host.hpp`

Mirror of device kernel, OpenMP-parallelized over n and y.

8.5 Python binding — `python/src/operators/py_op_brightness_contrast.cpp`

#include <pybind11/pybind11.h>
#include "op_brightness_contrast.hpp"
#include "py_tensor.hpp"

namespace py = pybind11;

void bind_brightness_contrast(py::module_& m) {
    m.def("brightness_contrast",
          [](py::object py_input,
             py::object py_output,
             py::object py_alpha,
             py::object py_beta,
             roccv::eDeviceType device,
             py::object py_stream) {
              auto in  = tensor_from_dlpack(py_input);
              auto out = tensor_from_dlpack(py_output);
              auto a   = tensor_from_dlpack(py_alpha);
              auto b   = tensor_from_dlpack(py_beta);
              hipStream_t stream = stream_from_pyobj(py_stream);
              roccv::BrightnessContrast op;
              op(stream, in, out, a, b, device);
          },
          py::arg("input"),
          py::arg("output"),
          py::arg("alpha"),
          py::arg("beta"),
          py::arg("device") = roccv::eDeviceType::GPU,
          py::arg("stream") = py::none(),
          "Apply per-image brightness/contrast adjustment.");
}

Register in python/src/main.cpp and add a stub to rocpycv.pyi.

8.6 Unit test — `tests/pybind/test_op_brightness_contrast.py`

import numpy as np
import pytest
import rocpycv as rcv
import cv2  # reference

@pytest.mark.parametrize("layout", ["NHWC", "HWC"])
@pytest.mark.parametrize("dtype", [np.uint8, np.float32])
@pytest.mark.parametrize("channels", [1, 3, 4])
@pytest.mark.parametrize("size", [(640, 480), (1920, 1080)])
def test_brightness_contrast_matches_opencv(layout, dtype, channels, size):
    N = 2 if layout == "NHWC" else 1
    H, W = size[1], size[0]
    shape = (N, H, W, channels) if layout == "NHWC" else (H, W, channels)
    rng = np.random.default_rng(0)
    src = rng.integers(0, 255, shape).astype(dtype) if dtype == np.uint8 \
          else rng.random(shape).astype(np.float32)
    alpha = np.array([1.4] * (N if layout == "NHWC" else 1), dtype=np.float32)
    beta  = np.array([10.0] * (N if layout == "NHWC" else 1), dtype=np.float32)

    rcv_in  = rcv.from_dlpack(src).copy_to(rcv.eDeviceType.GPU)
    rcv_out = rcv.empty_like(rcv_in)
    rcv_a   = rcv.from_dlpack(alpha).copy_to(rcv.eDeviceType.GPU)
    rcv_b   = rcv.from_dlpack(beta).copy_to(rcv.eDeviceType.GPU)
    rcv.brightness_contrast(rcv_in, rcv_out, rcv_a, rcv_b)
    got = rcv_out.copy_to(rcv.eDeviceType.CPU).numpy()

    if layout == "NHWC":
        ref = np.stack([cv2.convertScaleAbs(src[i], alpha=1.4, beta=10.0) for i in range(N)])
    else:
        ref = cv2.convertScaleAbs(src, alpha=1.4, beta=10.0)

    tol = 1 if dtype == np.uint8 else 1e-5
    assert np.max(np.abs(got.astype(int) - ref.astype(int))) <= tol

8.7 Benchmark — `benchmarks/src/bench_op_brightness_contrast.cpp`

Use the existing roccvbench harness; one row per (size, dtype, layout) reporting Mpix/s and µs/iter.

8.8 Docs — `docs/reference/operators/brightness_contrast.rst`

Standard operator page; copy flip.rst and adapt. Add a row to rocCV-supported-operators.rst.

9. Testing & validation strategy

9.1 Test matrix (per operator, minimum)

Axis	Values
Datatype	U8, F32 (and others if the op spec demands)
Layout	NHWC, HWC
Channels	1, 3, 4
Size	(640×480), (1920×1080); add (3840×2160) for the final perf sweep
Batch	1, 8
Device	CPU, GPU

That's at least 2 × 2 × 3 × 2 × 2 × 2 = 96 cases per op; cover with pytest.mark.parametrize and prune obviously redundant combinations.

9.2 Golden data

For each operator, generate reference output with the canonical library (in priority order):

torchvision.transforms.v2.functional if it exists (Pad, RandomResizedCrop, GaussianBlur).
opencv-python for everything else (Median, Adaptive, Morphology, HistEq, CLAHE).
Hand-derived analytic for the trivial ones (BrightnessContrast).

9.3 Tolerance

Datatype	Allowed max abs error vs reference
U8	≤ 1 (off-by-one rounding)
F32	≤ 1e-5 absolute or 1e-4 relative

If a particular op (HQResize, CLAHE) needs looser tolerance, document the reason in the PR description.

9.4 Edge cases test

Empty tensor (zero N or zero H/W) — should error or no-op cleanly.
Single-row / single-column image (W=1 or H=1).
Maximum supported dimensions (publish the spec, e.g. up to 8K × 8K).
Mismatched input/output shape — should throw roccv::Exception.
Wrong dtype / layout / channels — should throw with a clear message.
CPU vs GPU bit-for-bit (or within tolerance) parity.

10. Performance methodology

10.1 The benchmark harness

rocCV already has benchmarks/roccvbench/. Pattern: add a bench_op_<snake>.cpp that registers cases by (dtype, layout, channels, size, batch) and emits a JSON row consumable by analyze_results.py.

10.2 Required perf table per PR

For each operator, the PR description must include:

Size	rocCV GPU (Mpix/s)	rocCV CPU (Mpix/s)	OpenCV CPU (Mpix/s)	CV-CUDA GPU (if available, Mpix/s)
1920×1080 U8 3ch	…	…	…	…
3840×2160 U8 3ch	…	…	…	…
1920×1080 F32 3ch	…	…	…	…

10.3 Benchmark

Hardware: MI300X (or whatever AMD provides) on one node.
Workload: 50K-image subset, batch 256, image size 224×224.
Pipelines compared:
1. CPU torchvision.transforms.v2 (baseline).
2. rocCV-on-GPU via direct rocpycv calls (the demo).
3. (Optional) CV-CUDA pipeline on a comparable NVIDIA box if accessible.
Metric: images/second sustained for 100 steps, p50/p95 of per-step time.

11. References

Internal / upstream:

rocCV repo: https://github.com/ROCm/rocCV
rocCV docs: https://rocm.docs.amd.com/projects/rocCV/en/latest/
rocCV supported operators (stale, v0.2.0): https://rocm.docs.amd.com/projects/rocCV/en/latest/reference/rocCV-supported-operators.html
rocCV C++ how-to: https://rocm.docs.amd.com/projects/rocCV/en/latest/how-to/using-rocCV-cpp.html
rocCV Python how-to: https://rocm.docs.amd.com/projects/rocCV/en/latest/how-to/using-rocCV-python.html
RPP (sister library, for kernel patterns): https://github.com/ROCm/rpp

External references:

CV-CUDA operators: https://cvcuda.github.io/CV-CUDA/modules/python/operators.html
CV-CUDA developer guide: https://github.com/CVCUDA/CV-CUDA/blob/main/DEVELOPER_GUIDE.md
OpenCV image proc module: https://docs.opencv.org/4.x/d7/dbd/group__imgproc.html
torchvision Transforms v2: https://pytorch.org/vision/stable/transforms.html
DLPack spec: https://dmlc.github.io/dlpack/latest/

12. Stretch goals

Tier A — Architectural / standout impact

A1. Prototype `torchvision.transforms.v2` ROCm dispatch backend

A2. fp16 / bf16 datatype support across new operators

A3. Fused decode → resize → normalize pipeline (rocJPEG ↔ rocCV)

A4. Stand up a rocCV CI perf-gate (regression detection)

Tier B — Additional operators (extend the catalog)

#	Op	Effort	Reuse from main project	Why
B1	Sobel	2–3 days	GaussianBlur (separable conv)	Edge-detection staple; in every classical-CV pipeline.
B2	Laplacian	2 days	Sobel kernel pattern	Edge / blur-detection; pairs with Sobel.
B3	Conv2D (general)	5–7 days	GaussianBlur (uses smaller version of same kernel)	First "user supplies a kernel" op; opens door to lightweight feature extraction.
B4	Connected-components labeling	7–10 days	Histogram (reduction patterns)	Segmentation building block; first non-trivial GPU parallelism (union-find on device).
B5	OSD (on-screen display)	4–5 days	BndBox (already in develop)	Text / line / box overlay for inference demos; visible in every inference video.
B6	ChannelReorder / Stack	2 days	Reformat (already in develop)	Common preprocessing step; tiny but referenced everywhere.
B7	MorphologyEx with arbitrary structuring element	5–7 days	MorphologyEx (built earlier)	Generalizes beyond the 5 default modes; OpenCV-parity.
B8	HOG cells / HOG features	7–10 days	Histogram, Sobel	Classical detection; in the OpenVX 1.3 standard set, so doubles as a MIVisionX-side win.
B9	Inpaint (Telea / Navier-Stokes)	10+ days	None directly	Niche but CV-CUDA has it; image restoration / object removal.
B10	MinAreaRect + FindContours	7–10 days	NonMaximumSuppression patterns	Segmentation post-processing; CV-CUDA-parity.

Tier C — Ecosystem / developer experience

C1. Public docs refresh (`rocCV-supported-operators` page)

C2. HuggingFace `datasets` integration sample

C3. PyPI wheel for `rocpycv`

C4. DLPack v1 (managed-tensor versioned) support

C5. Multi-GPU pipeline sample

Tier D — Engineering hygiene

D1. Fuzz tests for tensor ops

D2. AddressSanitizer + UBSan CI pass

D3. clang-tidy / clang-format enforcement in CI

D4. rocprof integration sample

D5. Bench-result dashboard on GitHub Pages

Home
Contributing
Architecture
- Overview
- Containers
  - ImageBatchVarShape
  - Tensor
- Kernel Wrappers
- Operator Flow

Uh oh!

rocCV Operator Expansion

Goal

1. Summary

2. Background

2.1 What rocCV is

2.2 What's already in rocCV develop (21 operators)

2.3 The gap vs CV-CUDA

3. Problem statement

4. Goals and non-goals

4.1 Goals (must-have)

4.2 Non-goals (explicitly out of scope)

5. Definition of done

6. Target operator set

7. Codebase tour

8. The operator implementation template

8.1 Public header — include/op_brightness_contrast.hpp

8.2 Dispatch — src/op_brightness_contrast.cpp

8.3 HIP device kernel — include/kernels/device/brightness_contrast_device.hpp

8.4 Host fallback — include/kernels/host/brightness_contrast_host.hpp

8.5 Python binding — python/src/operators/py_op_brightness_contrast.cpp

8.6 Unit test — tests/pybind/test_op_brightness_contrast.py

8.7 Benchmark — benchmarks/src/bench_op_brightness_contrast.cpp

8.8 Docs — docs/reference/operators/brightness_contrast.rst

9. Testing & validation strategy

9.1 Test matrix (per operator, minimum)

9.2 Golden data

9.3 Tolerance

9.4 Edge cases test

10. Performance methodology

10.1 The benchmark harness

10.2 Required perf table per PR

10.3 Benchmark

11. References

12. Stretch goals

Tier A — Architectural / standout impact

A1. Prototype torchvision.transforms.v2 ROCm dispatch backend

A2. fp16 / bf16 datatype support across new operators

A3. Fused decode → resize → normalize pipeline (rocJPEG ↔ rocCV)

A4. Stand up a rocCV CI perf-gate (regression detection)

Tier B — Additional operators (extend the catalog)

Tier C — Ecosystem / developer experience

C1. Public docs refresh (rocCV-supported-operators page)

C2. HuggingFace datasets integration sample

C3. PyPI wheel for rocpycv

C4. DLPack v1 (managed-tensor versioned) support

C5. Multi-GPU pipeline sample

Tier D — Engineering hygiene

D1. Fuzz tests for tensor ops

D2. AddressSanitizer + UBSan CI pass

D3. clang-tidy / clang-format enforcement in CI

D4. rocprof integration sample

D5. Bench-result dashboard on GitHub Pages

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Clone this wiki locally

2.2 What's already in `rocCV develop` (21 operators)

8.1 Public header — `include/op_brightness_contrast.hpp`

8.2 Dispatch — `src/op_brightness_contrast.cpp`

8.3 HIP device kernel — `include/kernels/device/brightness_contrast_device.hpp`

8.4 Host fallback — `include/kernels/host/brightness_contrast_host.hpp`

8.5 Python binding — `python/src/operators/py_op_brightness_contrast.cpp`

8.6 Unit test — `tests/pybind/test_op_brightness_contrast.py`

8.7 Benchmark — `benchmarks/src/bench_op_brightness_contrast.cpp`

8.8 Docs — `docs/reference/operators/brightness_contrast.rst`

A1. Prototype `torchvision.transforms.v2` ROCm dispatch backend

C1. Public docs refresh (`rocCV-supported-operators` page)

C2. HuggingFace `datasets` integration sample

C3. PyPI wheel for `rocpycv`