[New Operator] Initial Implementation of AverageBlur

CHANGELOG.md

@@ -8,6 +8,7 @@ The full documentation for rocCV is available at [https://rocm.docs.amd.com/proj
 - Added `Reformat` operator.
 - Added `BrightnessContrast` operator.
 - Added `Gaussian` operator.
+- Added `AverageBlur` operator.
 - Dockerfile support for running rocCV in a container environment.
 
 ### Changed

benchmarks/src/roccv/bench_average_blur.cpp

@@ -0,0 +1,123 @@
+/*
+ * 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.
+ */
+
+#include <core/hip_assert.h>
+
+#include <cfenv>
+
+#include <core/image_format.hpp>
+#include <core/tensor.hpp>
+#include <op_average_blur.hpp>
+#include <roccvbench/registry.hpp>
+#include <roccvbench/utils.hpp>
+
+#include "roccv_bench_helpers.hpp"
+
+using namespace roccv;
+
+template <eDeviceType DeviceType>
+static roccvbench::BenchmarkResults RunAverageBlurBenchmark(roccvbench::BenchmarkParamsList params) {
+    roccvbench::BenchmarkResults results;
+
+    int samples = roccvbench::GetParamValue<int>(params, "samples");
+    int width = roccvbench::GetParamValue<int>(params, "width");
+    int height = roccvbench::GetParamValue<int>(params, "height");
+    int runs = roccvbench::GetParamValue<int>(params, "runs");
+    int warmupRuns = roccvbench::GetParamValue<int>(params, "warmupRuns");
+    ImageFormat inFormat = roccvbench::GetParamValue<ImageFormat>(params, "inFormat");
+    ImageFormat outFormat = roccvbench::GetParamValue<ImageFormat>(params, "outFormat");
+    int kernelWidth = roccvbench::GetParamValue<int>(params, "kernelWidth");
+    int kernelHeight = roccvbench::GetParamValue<int>(params, "kernelHeight");
+    eBorderType borderType = roccvbench::GetParamValue<eBorderType>(params, "borderType");
+
+    TensorRequirements inReqs = Tensor::CalcRequirements(samples, {width, height}, inFormat, DeviceType);
+    TensorRequirements outReqs = Tensor::CalcRequirements(samples, {width, height}, outFormat, DeviceType);
+    Tensor input(inReqs);
+    Tensor output(outReqs);
+
+    RegisterMemoryUsage(input, results.readMemoryBytes);
+    RegisterMemoryUsage(output, results.writtenMemoryBytes);
+
+    FillTensor(input);
+
+    AverageBlur op(kernelWidth, kernelHeight);
+
+    hipStream_t stream;
+    HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream));
+
+    roccvbench::RecordRuns<DeviceType>(stream, runs, warmupRuns, results.executionTimes, [&]() {
+        op(stream, input, output, kernelWidth, kernelHeight, -1, -1, borderType, DeviceType);
+    });
+
+    HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream));
+
+    return results;
+}
+
+#define DEFINE_AVERAGE_BLUR_BENCHMARK(name, device, inFormat, outFormat, kernelWidth, kernelHeight, borderType) \
+    BENCHMARK_P(AverageBlur, name,                                                                                  \
+                BENCH_PARAMS(BENCH_PARAM("inFormat", inFormat), BENCH_PARAM("outFormat", outFormat),             \
+                             BENCH_PARAM("kernelWidth", kernelWidth), BENCH_PARAM("kernelHeight", kernelHeight), \
+                             BENCH_PARAM("borderType", borderType))) {                                               \
+        return RunAverageBlurBenchmark<device>(params);                                                             \
+    }
+
+// GPU benchmarks
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 3, 3, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 3, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 3, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 3, 9, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 5, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 7, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 9, 9, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 11, 11, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 15, 15, eBorderType::BORDER_TYPE_REFLECT);
+
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 3, 3, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 3, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 3, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 3, 9, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 5, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 7, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 9, 9, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 11, 11, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 15, 15, eBorderType::BORDER_TYPE_REFLECT);
+
+// CPU benchmarks
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 3, 3, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 3, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 3, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 3, 9, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 5, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 7, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 9, 9, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 11, 11, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 15, 15, eBorderType::BORDER_TYPE_REFLECT);
+
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 3, 3, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 3, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 3, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 3, 9, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 5, 5, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 7, 7, eBorderType::BORDER_TYPE_REFLECT);
+DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 9, 9, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 11, 11, eBorderType::BORDER_TYPE_REFLECT);
+//DEFINE_AVERAGE_BLUR_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 15, 15, eBorderType::BORDER_TYPE_REFLECT);

docs/reference/rocCV-supported-operators.rst

@@ -13,6 +13,7 @@ The rocCV is a collection of the following computer vision operators that are su
     :header: "Operator", "Description", "Datatypes", "Layouts"
 
     "AdvCvtColor","Converts color spaces with explicit BT601/BT709/BT2020 coefficients, including NV12/NV21 paths.","U8","NHWC, HWC"
+    "AverageBlur","Applies an average filter.","U8, U16, S16, S32, F32", "NHWC, HWC"
     "BilateralFilter", "Applies a bilateral filter.", "U8", "NHWC, HWC"
     "BndBox","Draws bounding boxes on the images in a tensor.","U8","NHWC, HWC"
     "BrightnessContrast", "Adjusts the brightness and contrast on images in a tensor.", "U8, U16, S16, S32, F32", "NHWC, HWC"

docs/supported-operators.md

@@ -5,6 +5,7 @@ See below for a list of Computer Vision operators rocCV supports.
 |Name|Description|Datatypes|Layouts|CPU/GPU Support|
 |-|-|-|-|-|
 |AdvCvtColor|Converts color spaces using explicit BT601/BT709/BT2020 coefficients and supports NV12/NV21 paths.|U8|NHWC, HWC|Both|
+|AverageBlur|Applies an average filter for image smoothing.|U8, U16, S16, S32, F32|NHWC, HWC|Both|
 |BilateralFilter|Applies a bilateral filter to reduce image noise while preserving strong edges.|U8, S8, U16, S16, U32, S32, F32, F64|NHWC, HWC|Both|
 |BndBox|Draws rectangular borders using the specified locations, dimensions and colors, in order to show the locations and sizes of objects in an image.|U8, S8|NHWC, HWC|Both|
 |BrightnessContrast|Adjusts the brightness and contrast of an image.|U8, U16, S16, S32, F32|NHWC, HWC|Both|

include/op_average_blur.hpp

@@ -0,0 +1,113 @@
+/**
+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 <mutex>
+
+#include "core/detail/allocators/default_allocator.hpp"
+#include "core/tensor.hpp"
+#include "i_operator.hpp"
+#include "operator_types.h"
+
+namespace roccv {
+/**
+ * @brief Class for managing the Average Blur operator.
+ *
+ */
+class AverageBlur : public IOperator {
+   public:
+    /**
+     * @brief Constructs an AverageBlur object.
+     *
+     * @param maxKernelWidth [in] The maximum kernel width that will be used by the operator. Must be positive (> 0).
+     * @param maxKernelHeight [in] The maximum kernel height that will be used by the operator. Must be positive(> 0).
+     */
+    AverageBlur(int32_t maxKernelWidth, int32_t maxKernelHeight);
+
+    /**
+     * @brief Destroy the AverageBlur object
+     *
+     */
+    ~AverageBlur();
+
+    /**
+     * @brief Executes the Average Blur operation on the given HIP stream.
+     *
+     *
+     * Limitations:
+     *
+     * Input:
+     *       Supported TensorLayout(s): [NHWC, HWC]
+     *                        Channels: [1, 3, 4]
+     *       Supported DataType(s):     [U8, U16, S16, S32, F32]
+     *
+     * Output:
+     *       Supported TensorLayout(s): [NHWC, HWC]
+     *                        Channels: [1, 3, 4]
+     *       Supported DataType(s):     [U8, U16, S16, S32, F32]
+     *
+     *
+     * Input/Output dependency:
+     *
+     *       Property      |  Input == Output
+     *      -------------- | -------------
+     *       TensorLayout  | Yes
+     *       DataType      | Yes
+     *       Channels      | Yes
+     *       Width         | Yes
+     *       Height        | Yes
+     *       Batch         | Yes
+     *
+     * @param[in] stream The HIP stream to run this operation on.
+     * @param[in] input Input tensor.
+     * @param[out] output Output tensor.
+     * @param[in] kernelWidth AverageBlur kernel width. Must be positive, odd, and less than m_maxKernelWidth.
+     * @param[in] kernelHeight AverageBlur kernel height. Must be positive, odd, and less than m_maxKernelHeight.
+     * @param[in] anchorX Kernel anchor in X direction. Must be nonnegative and less than kernelWidth, or -1 to indicate
+     * center.
+     * @param[in] anchorY Kernel anchor in Y direction. Must be nonnegative and less than kernelHeight, or -1 to
+     * indicate center.
+     * @param[in] borderMode A border type to identify the pixel extrapolation
+     * method (e.g. BORDER_TYPE_CONSTANT or BORDER_TYPE_REPLICATE).
+     * @param[in] device The device which this operation should run on. (Default: eDeviceType::GPU)
+     */
+    void operator()(hipStream_t stream, const Tensor& input, Tensor& output, int kernelWidth, int kernelHeight,
+                    int anchorX, int anchorY, eBorderType borderMode, eDeviceType device = eDeviceType::GPU);
+
+   private:
+    int32_t m_maxKernelWidth;
+    int32_t m_maxKernelHeight;
+    DefaultAllocator m_allocator;
+    std::mutex m_bufferMutex;
+    bool m_use2D = false;
+    float* m_hostKernelMem = nullptr;
+    float* m_hostKernelMemH = nullptr;
+    float* m_hostKernelMemV = nullptr;
+    float* m_deviceKernelMem = nullptr;
+    float* m_deviceKernelMemH = nullptr;
+    float* m_deviceKernelMemV = nullptr;
+    hipEvent_t m_completionEvent = nullptr;
+};
+}  // namespace roccv

include/roccv_operators.hpp

@@ -21,6 +21,7 @@ THE SOFTWARE.
 */
 
 #include "op_adv_cvt_color.hpp"
+#include "op_average_blur.hpp"
 #include "op_bilateral_filter.hpp"
 #include "op_bnd_box.hpp"
 #include "op_brightness_contrast.hpp"

python/include/operators/py_op_average_blur.hpp

@@ -0,0 +1,42 @@
+/**
+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 <operator_types.h>
+#include <pybind11/pybind11.h>
+
+#include "py_stream.hpp"
+#include "py_tensor.hpp"
+
+namespace py = pybind11;
+
+class PyOpAverageBlur {
+   public:
+    static void Export(py::module& m);
+    static PyTensor Execute(PyTensor& input, const std::tuple<int, int>& kernelSize,
+                            const std::tuple<int, int>& kernelAnchor, eBorderType borderMode,
+                            std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device);
+    static void ExecuteInto(PyTensor& output, PyTensor& input, const std::tuple<int, int>& kernelSize,
+                            const std::tuple<int, int>& kernelAnchor, eBorderType borderMode,
+                            std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device);
+};

python/src/main.cpp

@@ -24,6 +24,7 @@ THE SOFTWARE.
 #include <pybind11/stl.h>
 
 #include "operators/py_op_adv_cvt_color.hpp"
+#include "operators/py_op_average_blur.hpp"
 #include "operators/py_op_bilateral_filter.hpp"
 #include "operators/py_op_bnd_box.hpp"
 #include "operators/py_op_brightness_contrast.hpp"
@@ -63,6 +64,7 @@ PYBIND11_MODULE(rocpycv, m) {
     PyStructs::Export(m);
     PyStream::Export(m);
     PyTensor::Export(m);
+    PyOpAverageBlur::Export(m);
     PyOpCustomCrop::Export(m);
     PyOpNonMaxSuppression::Export(m);
     PyOpNormalize::Export(m);