diff --git a/CHANGELOG.md b/CHANGELOG.md index 245c5fca..ff45bb01 100644 --- a/CHANGELOG.md +++ b/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 diff --git a/benchmarks/src/roccv/bench_average_blur.cpp b/benchmarks/src/roccv/bench_average_blur.cpp new file mode 100644 index 00000000..442da87c --- /dev/null +++ b/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 + +#include + +#include +#include +#include +#include +#include + +#include "roccv_bench_helpers.hpp" + +using namespace roccv; + +template +static roccvbench::BenchmarkResults RunAverageBlurBenchmark(roccvbench::BenchmarkParamsList params) { + roccvbench::BenchmarkResults results; + + int samples = roccvbench::GetParamValue(params, "samples"); + int width = roccvbench::GetParamValue(params, "width"); + int height = roccvbench::GetParamValue(params, "height"); + int runs = roccvbench::GetParamValue(params, "runs"); + int warmupRuns = roccvbench::GetParamValue(params, "warmupRuns"); + ImageFormat inFormat = roccvbench::GetParamValue(params, "inFormat"); + ImageFormat outFormat = roccvbench::GetParamValue(params, "outFormat"); + int kernelWidth = roccvbench::GetParamValue(params, "kernelWidth"); + int kernelHeight = roccvbench::GetParamValue(params, "kernelHeight"); + eBorderType borderType = roccvbench::GetParamValue(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(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(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); \ No newline at end of file diff --git a/docs/reference/rocCV-supported-operators.rst b/docs/reference/rocCV-supported-operators.rst index 2d80df3d..69779a95 100644 --- a/docs/reference/rocCV-supported-operators.rst +++ b/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" diff --git a/docs/supported-operators.md b/docs/supported-operators.md index 78ba9b86..07945f46 100644 --- a/docs/supported-operators.md +++ b/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| diff --git a/include/op_average_blur.hpp b/include/op_average_blur.hpp new file mode 100644 index 00000000..0010cab7 --- /dev/null +++ b/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 + +#include + +#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 diff --git a/include/roccv_operators.hpp b/include/roccv_operators.hpp index d3791de7..c78a7a3e 100644 --- a/include/roccv_operators.hpp +++ b/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" diff --git a/python/include/operators/py_op_average_blur.hpp b/python/include/operators/py_op_average_blur.hpp new file mode 100644 index 00000000..31154c5c --- /dev/null +++ b/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 +#include + +#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& kernelSize, + const std::tuple& kernelAnchor, eBorderType borderMode, + std::optional> stream, eDeviceType device); + static void ExecuteInto(PyTensor& output, PyTensor& input, const std::tuple& kernelSize, + const std::tuple& kernelAnchor, eBorderType borderMode, + std::optional> stream, eDeviceType device); +}; \ No newline at end of file diff --git a/python/src/main.cpp b/python/src/main.cpp index 82c7ffee..cf64473f 100644 --- a/python/src/main.cpp +++ b/python/src/main.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. #include #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); diff --git a/python/src/operators/py_op_average_blur.cpp b/python/src/operators/py_op_average_blur.cpp new file mode 100644 index 00000000..a5a88162 --- /dev/null +++ b/python/src/operators/py_op_average_blur.cpp @@ -0,0 +1,105 @@ +/** +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 "operators/py_op_average_blur.hpp" + +#include + +#include "py_helpers.hpp" + +PyTensor PyOpAverageBlur::Execute(PyTensor& input, const std::tuple& kernelSize, + const std::tuple& kernelAnchor, eBorderType borderMode, + std::optional> stream, eDeviceType device) { + hipStream_t hipStream = stream.has_value() ? stream.value().get().getStream() : nullptr; + + auto inputTensor = input.getTensor(); + auto outputTensor = std::make_shared(inputTensor->shape(), inputTensor->dtype(), device); + + int kernelWidth = std::get<0>(kernelSize); + int kernelHeight = std::get<1>(kernelSize); + + roccv::AverageBlur op(kernelWidth, kernelHeight); + op(hipStream, *inputTensor, *outputTensor, kernelWidth, kernelHeight, std::get<0>(kernelAnchor), + std::get<1>(kernelAnchor), borderMode, device); + return PyTensor(outputTensor); +} + +void PyOpAverageBlur::ExecuteInto(PyTensor& output, PyTensor& input, const std::tuple& kernelSize, + const std::tuple& kernelAnchor, eBorderType borderMode, + std::optional> stream, eDeviceType device) { + hipStream_t hipStream = stream.has_value() ? stream.value().get().getStream() : nullptr; + int kernelWidth = std::get<0>(kernelSize); + int kernelHeight = std::get<1>(kernelSize); + + roccv::AverageBlur op(kernelWidth, kernelHeight); + op(hipStream, *input.getTensor(), *output.getTensor(), kernelWidth, kernelHeight, std::get<0>(kernelAnchor), + std::get<1>(kernelAnchor), borderMode, device); +} + +void PyOpAverageBlur::Export(py::module& m) { + using namespace py::literals; + const std::tuple default_anchor{-1, -1}; + + m.def("averageblur", &PyOpAverageBlur::Execute, "src"_a, "kernelSize"_a, "kernelAnchor"_a = default_anchor, + "borderMode"_a = BORDER_TYPE_CONSTANT, py::kw_only(), "stream"_a = nullptr, "device"_a = eDeviceType::GPU, + R"pbdoc( + + Executes the AverageBlur operation on the given HIP stream. + + See also: + Refer to the rocCV C++ API reference for more information on this operation. + + Args: + src (rocpycv.Tensor): Input tensor containing one or more images. + kernelSize (Tuple[int, int]): AverageBlur kernel width, height. Both kernel width and height must be odd and positive. + kernelAnchor (Tuple[int, int]): AverageBlur kernel anchor in X,Y directions. Must be positive and less than kernelSize, or -1 to indicate kernel center. Defaults to (-1, -1). + borderMode (rocpycv.eBorderType, optional): The border type to identify the pixel extrapolation method. Defaults to BORDER_TYPE_CONSTANT. + stream (rocpycv.Stream, optional): HIP stream to run this operation on. + device (rocpycv.Device, optional): The device to run this operation on. Defaults to GPU. + + Returns: + rocpycv.Tensor: The output tensor. + )pbdoc"); + + m.def("averageblur_into", &PyOpAverageBlur::ExecuteInto, "dst"_a, "src"_a, "kernelSize"_a, + "kernelAnchor"_a = default_anchor, "borderMode"_a = BORDER_TYPE_CONSTANT, py::kw_only(), "stream"_a = nullptr, + "device"_a = eDeviceType::GPU, + R"pbdoc( + + Executes the AverageBlur operation on the given HIP stream. + + See also: + Refer to the rocCV C++ API reference for more information on this operation. + + Args: + dst (rocpycv.Tensor): The output tensor which results are written to. + src (rocpycv.Tensor): Input tensor containing one or more images. + kernelSize (Tuple[int, int]): AverageBlur kernel width, height. Both kernel width and height must be odd and positive. + kernelAnchor (Tuple[int, int]): AverageBlur kernel anchor in X,Y directions. Must each be positive and less than kernelSize, or -1 to indicate kernel center. Defaults to (-1, -1). + borderMode (rocpycv.eBorderType, optional): The border type to identify the pixel extrapolation method. Defaults to BORDER_TYPE_CONSTANT. + stream (rocpycv.Stream, optional): HIP stream to run this operation on. + device (rocpycv.Device, optional): The device to run this operation on. Defaults to GPU. + + Returns: + None + )pbdoc"); +} \ No newline at end of file diff --git a/src/op_average_blur.cpp b/src/op_average_blur.cpp new file mode 100644 index 00000000..8e3f1060 --- /dev/null +++ b/src/op_average_blur.cpp @@ -0,0 +1,224 @@ +/** +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 "op_average_blur.hpp" + +#include + +#include +#include +#include + +#include "common/validation_helpers.hpp" +#include "core/detail/casting.hpp" +#include "details/filter_2d.hpp" + +namespace roccv { +AverageBlur::AverageBlur(int32_t maxKernelWidth, int32_t maxKernelHeight) + : m_maxKernelWidth(maxKernelWidth), m_maxKernelHeight(maxKernelHeight) { + if (maxKernelWidth <= 0) { + throw roccv::Exception( + "Invalid maxKernelWidth = " + std::to_string(maxKernelWidth) + ": Ensure that it is positive.", + eStatusType::INVALID_VALUE); + } + if (maxKernelHeight <= 0) { + throw roccv::Exception( + "Invalid maxKernelHeight = " + std::to_string(maxKernelHeight) + ": Ensure that it is positive.", + eStatusType::INVALID_VALUE); + } + + size_t memSize = maxKernelHeight * maxKernelWidth * sizeof(float); + size_t memSizeH = m_maxKernelWidth * sizeof(float); + size_t memSizeV = m_maxKernelHeight * sizeof(float); + + if (maxKernelHeight <= 3 && maxKernelWidth <= 3) { + m_use2D = true; + m_hostKernelMem = static_cast(m_allocator.allocHostPinnedMem(memSize)); + } else { + m_hostKernelMemH = static_cast(m_allocator.allocHostPinnedMem(memSizeH)); + m_hostKernelMemV = static_cast(m_allocator.allocHostPinnedMem(memSizeV)); + } + + int gpuCount = 0; + hipError_t err = hipGetDeviceCount(&gpuCount); + if (err == hipSuccess) { + if (m_use2D) { + m_deviceKernelMem = static_cast(m_allocator.allocHipMem(memSize)); + } else { + m_deviceKernelMemH = static_cast(m_allocator.allocHipMem(memSizeH)); + m_deviceKernelMemV = static_cast(m_allocator.allocHipMem(memSizeV)); + } + HIP_VALIDATE_NO_ERRORS(hipEventCreateWithFlags(&m_completionEvent, hipEventDisableTiming)); + } +} + +AverageBlur::~AverageBlur() { + if (m_use2D) { + m_allocator.freeHostPinnedMem(m_hostKernelMem); + } else { + m_allocator.freeHostPinnedMem(m_hostKernelMemH); + m_allocator.freeHostPinnedMem(m_hostKernelMemV); + } + if (m_deviceKernelMem != nullptr) { + m_allocator.freeHipMem(m_deviceKernelMem); + } + if (m_deviceKernelMemH != nullptr) { + m_allocator.freeHipMem(m_deviceKernelMemH); + } + if (m_deviceKernelMemV != nullptr) { + m_allocator.freeHipMem(m_deviceKernelMemV); + } + if (m_completionEvent != nullptr) { + (void)hipEventDestroy(m_completionEvent); + } +} + +void AverageBlur::operator()(hipStream_t stream, const Tensor& input, Tensor& output, int kernelWidth, int kernelHeight, + int anchorX, int anchorY, eBorderType borderMode, eDeviceType device) { + // Validate input tensor + CHECK_TENSOR_DEVICE(input, device); + CHECK_TENSOR_DATATYPES(input, DATA_TYPE_U8, DATA_TYPE_U16, DATA_TYPE_S16, DATA_TYPE_S32, DATA_TYPE_F32); + CHECK_TENSOR_LAYOUT(input, TENSOR_LAYOUT_HWC, TENSOR_LAYOUT_NHWC); + CHECK_TENSOR_CHANNELS(input, 1, 3, 4); + + // Validate output tensor + CHECK_TENSOR_COMPARISON(input.dtype() == output.dtype()); + CHECK_TENSOR_COMPARISON(input.device() == output.device()); + CHECK_TENSOR_COMPARISON(input.shape() == output.shape()); + + // Validate kernel size + if (!(kernelWidth > 0 && kernelWidth % 2 == 1 && kernelWidth <= m_maxKernelWidth && kernelHeight > 0 && + kernelHeight % 2 == 1 && kernelHeight <= m_maxKernelHeight)) { + throw roccv::Exception("Invalid kernel size = " + std::to_string(kernelWidth) + ", " + + std::to_string(kernelHeight) + + ": Ensure that the kernel size is odd, positive, and less than the max:" + + std::to_string(m_maxKernelWidth) + ", " + std::to_string(m_maxKernelHeight), + eStatusType::INVALID_VALUE); + } + + // Validate anchor + if (!((anchorX == -1 || (anchorX >= 0 && anchorX < kernelWidth)) && + (anchorY == -1 || (anchorY >= 0 && anchorY < kernelHeight)))) { + throw roccv::Exception("Invalid anchor = " + std::to_string(anchorX) + ", " + std::to_string(anchorY) + + ": Ensure that the anchorX and anchorY are -1, or nonnegative and less than " + "kernelWidth and kernelHeight, respectively:" + + std::to_string(kernelWidth) + ", " + std::to_string(kernelHeight), + eStatusType::INVALID_VALUE); + } + processAnchor(anchorX, anchorY, kernelWidth, kernelHeight); + + std::lock_guard lock(m_bufferMutex); + if (device == eDeviceType::GPU) { + HIP_VALIDATE_NO_ERRORS(hipEventSynchronize(m_completionEvent)); + } + + // Compute the kernel + if (m_use2D) { + float val = 1.0 / (kernelWidth * kernelHeight); + for (int y = 0; y < kernelHeight; ++y) { + for (int x = 0; x < kernelWidth; ++x) { + m_hostKernelMem[y * kernelWidth + x] = val; + } + } + } else { + float valX = 1.0 / kernelWidth; + for (int x = 0; x < kernelWidth; ++x) { + m_hostKernelMemH[x] = valX; + } + float valY = 1.0 / kernelHeight; + for (int y = 0; y < kernelHeight; ++y) { + m_hostKernelMemV[y] = valY; + } + } + + if (device == eDeviceType::GPU) { + if (m_use2D) { + if (m_deviceKernelMem == nullptr) { + throw roccv::Exception("Device memory not allocated for AverageBlur kernel, GPU may not be available.", + eStatusType::INVALID_OPERATION); + } + HIP_VALIDATE_NO_ERRORS(hipMemcpyAsync(m_deviceKernelMem, m_hostKernelMem, + kernelWidth * kernelHeight * sizeof(float), hipMemcpyHostToDevice, + stream)); + } else { + if (m_deviceKernelMemH == nullptr || m_deviceKernelMemV == nullptr) { + throw roccv::Exception("Device memory not allocated for AverageBlur kernel, GPU may not be available.", + eStatusType::INVALID_OPERATION); + } + HIP_VALIDATE_NO_ERRORS(hipMemcpyAsync(m_deviceKernelMemH, m_hostKernelMemH, kernelWidth * sizeof(float), + hipMemcpyHostToDevice, stream)); + HIP_VALIDATE_NO_ERRORS(hipMemcpyAsync(m_deviceKernelMemV, m_hostKernelMemV, kernelHeight * sizeof(float), + hipMemcpyHostToDevice, stream)); + } + } + + if (m_use2D) { + // clang-format off + static const std::unordered_map< + eDataType, std::array, 4>> + funcs = + { + {eDataType::DATA_TYPE_U8, {dispatch_filter_2d_dtype, 0, dispatch_filter_2d_dtype, dispatch_filter_2d_dtype}}, + {eDataType::DATA_TYPE_U16, {dispatch_filter_2d_dtype, 0, dispatch_filter_2d_dtype, dispatch_filter_2d_dtype}}, + {eDataType::DATA_TYPE_S16, {dispatch_filter_2d_dtype, 0, dispatch_filter_2d_dtype, dispatch_filter_2d_dtype}}, + {eDataType::DATA_TYPE_S32, {dispatch_filter_2d_dtype, 0, dispatch_filter_2d_dtype, dispatch_filter_2d_dtype}}, + {eDataType::DATA_TYPE_F32, {dispatch_filter_2d_dtype, 0, dispatch_filter_2d_dtype, dispatch_filter_2d_dtype}}, + }; + // clang-format on + auto func = funcs.at(input.dtype().etype())[input.shape(input.layout().channels_index()) - 1]; + if (func == 0) throw Exception("Not mapped to a defined function.", eStatusType::INVALID_OPERATION); + + if (device == eDeviceType::GPU) { + func(stream, input, output, m_deviceKernelMem, kernelWidth, kernelHeight, anchorX, anchorY, borderMode, + device); + HIP_VALIDATE_NO_ERRORS(hipEventRecord(m_completionEvent, stream)); + } else if (device == eDeviceType::CPU) { + func(stream, input, output, m_hostKernelMem, kernelWidth, kernelHeight, anchorX, anchorY, borderMode, + device); + } + } else { + // clang-format off + static const std::unordered_map< + eDataType, std::array, 4>> + funcs = + { + {eDataType::DATA_TYPE_U8, {dispatch_filter_2d_dtype_separable, 0, dispatch_filter_2d_dtype_separable, dispatch_filter_2d_dtype_separable}}, + {eDataType::DATA_TYPE_U16, {dispatch_filter_2d_dtype_separable, 0, dispatch_filter_2d_dtype_separable, dispatch_filter_2d_dtype_separable}}, + {eDataType::DATA_TYPE_S16, {dispatch_filter_2d_dtype_separable, 0, dispatch_filter_2d_dtype_separable, dispatch_filter_2d_dtype_separable}}, + {eDataType::DATA_TYPE_S32, {dispatch_filter_2d_dtype_separable, 0, dispatch_filter_2d_dtype_separable, dispatch_filter_2d_dtype_separable}}, + {eDataType::DATA_TYPE_F32, {dispatch_filter_2d_dtype_separable, 0, dispatch_filter_2d_dtype_separable, dispatch_filter_2d_dtype_separable}}, + }; + // clang-format on + auto func = funcs.at(input.dtype().etype())[input.shape(input.layout().channels_index()) - 1]; + if (func == 0) throw Exception("Not mapped to a defined function.", eStatusType::INVALID_OPERATION); + + if (device == eDeviceType::GPU) { + func(stream, input, output, m_deviceKernelMemH, m_deviceKernelMemV, kernelWidth, kernelHeight, anchorX, + anchorY, borderMode, device); + HIP_VALIDATE_NO_ERRORS(hipEventRecord(m_completionEvent, stream)); + } else if (device == eDeviceType::CPU) { + func(stream, input, output, m_hostKernelMemH, m_hostKernelMemV, kernelWidth, kernelHeight, anchorX, anchorY, + borderMode, device); + } + } +} +} // namespace roccv \ No newline at end of file diff --git a/tests/roccv/cpp/src/tests/operators/test_op_average_blur.cpp b/tests/roccv/cpp/src/tests/operators/test_op_average_blur.cpp new file mode 100644 index 00000000..b8b706d9 --- /dev/null +++ b/tests/roccv/cpp/src/tests/operators/test_op_average_blur.cpp @@ -0,0 +1,393 @@ +/** +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 +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "test_helpers.hpp" + +using namespace roccv; +using namespace roccv::detail; +using namespace roccv::tests; + +namespace { + +/** + * @brief Golden model for the AverageBlur operation. + * + * @tparam T Vectorized datatype of the image's pixels. + * @tparam BorderMode Border pixel extrapolation method. + * @tparam BT Base type of the image's data. + * @param[in] input Input tensor containing image data. + * @param[in] batchSize The number of images in the batch. + * @param[in] width Image width. + * @param[in] height Image height. + * @param[in] kernelWidth Kernel width. + * @param[in] kernelHeight Kernel height. + * @param[in] anchorX Kernel anchor in X direction. + * @param[in] anchorY Kernel anchor in Y direction. + * @return Vector containing the results of the operation. + */ +template > +std::vector GenerateGoldenAverageBlur(std::vector& input, int32_t batchSize, int32_t width, int32_t height, + int kernelWidth, int kernelHeight, int anchorX, int anchorY) { + std::vector output(input.size()); + BorderWrapper src(ImageWrapper(input, batchSize, width, height), SetAll(0)); + ImageWrapper dst(output, batchSize, width, height); + + using namespace roccv::detail; + using worktype = MakeType>; + + for (int b = 0; b < dst.batches(); b++) { + for (int j = 0; j < dst.height(); j++) { + for (int i = 0; i < dst.width(); i++) { + worktype numerators = SetAll(0.0f); + float denominator = 0.0f; + + for (int y = j - anchorY; y <= j + (kernelHeight - 1 - anchorY); y++) { + for (int x = i - anchorX; x <= i + (kernelWidth - 1 - anchorX); x++) { + worktype workPixel = StaticCast(src.at(b, y, x, 0)); + + denominator += 1.0f; + numerators += workPixel; + } + } + dst.at(b, j, i, 0) = SaturateCast(numerators / denominator); + } + } + } + return output; +} + +/** + * @brief Tests correctness of the AverageBlur operator, comparing it against a generated golden result. + * + * @tparam T Underlying datatype of the image's pixels. + * @tparam BorderMode Border pixel extrapolation method. + * @tparam BT Base type of the image's data. + * @param[in] batchSize Number of images in the batch. + * @param[in] width Width of each image in the batch. + * @param[in] height Height of each image in the batch. + * @param[in] format Image format. + * @param[in] kernelWidth Kernel width. + * @param[in] kernelHeight Kernel height. + * @param[in] anchorX Kernel anchor in X direction. + * @param[in] anchorY Kernel anchor in Y direction. + * @param[in] device Device this correctness test should be run on. + */ +template > +void TestCorrectness(int batchSize, int width, int height, ImageFormat format, int kernelWidth, int kernelHeight, + int anchorX, int anchorY, eDeviceType device) { + // Create input and output tensor based on test parameters + Tensor input(batchSize, {width, height}, format, device); + Tensor output(batchSize, {width, height}, format, device); + + // Create a vector and fill it with random data. + std::vector inputData(input.shape().size()); + FillVector(inputData); + if constexpr (std::is_floating_point_v) { + for (size_t i = 0; i < inputData.size(); i++) { + inputData[i] *= static_cast(std::numeric_limits::max()); + } + } + + // Copy generated input data into input tensor + CopyVectorIntoTensor(input, inputData); + + // Infer anchor if (-1,-1) + int effectiveAnchorX = (anchorX == -1) ? kernelWidth >> 1 : anchorX; + int effectiveAnchorY = (anchorY == -1) ? kernelHeight >> 1 : anchorY; + + hipStream_t stream; + HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream)); + AverageBlur op(kernelWidth, kernelHeight); + op(stream, input, output, kernelWidth, kernelHeight, anchorX, anchorY, BorderMode, device); + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(stream)); + HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream)); + + // Copy data from output tensor into a host allocated vector + std::vector outputData(output.shape().size()); + CopyTensorIntoVector(outputData, output); + + // Calculate golden reference + std::vector ref = GenerateGoldenAverageBlur(inputData, batchSize, width, height, kernelWidth, + kernelHeight, effectiveAnchorX, effectiveAnchorY); + + // Compare data in actual output versus the generated golden reference image + // TODO check on delta, this matches bilateral filter and passes (looser does not) + CompareVectorsNear(outputData, ref, 1); +} + +/** + * @brief Tests correctness of the AverageBlur operator when multiple threads concurrently use the same operator. + * @tparam T Underlying datatype of the image's pixels. + * @tparam BorderMode Border pixel extrapolation method. + * @tparam BT Base type of the image's data. + * @param[in] batchSize Number of images in the batch. + * @param[in] width Width of each image in the batch. + * @param[in] height Height of each image in the batch. + * @param[in] format Image format. + * @param[in] device Device this correctness test should be run on. + */ +template > +void TestCorrectnessConcurrent(int batchSize, int width, int height, ImageFormat format, eDeviceType device) { + constexpr int NUM_THREADS = 8; + constexpr int TOTAL_TESTS = 80; + + struct ThreadTest { + Tensor input; + Tensor output; + std::vector inputData; + hipStream_t stream; + int ksize; + + ThreadTest(int b, int w, int h, ImageFormat fmt, eDeviceType dev, int k, int id) + : input(b, {w, h}, fmt, dev), output(b, {w, h}, fmt, dev), inputData(input.shape().size()), ksize(k) { + HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream)); + FillVector(inputData, id * 1000); + CopyVectorIntoTensor(input, inputData); + } + + ~ThreadTest() { (void)hipStreamDestroy(stream); } + }; + + std::vector> threadTests; + threadTests.reserve(TOTAL_TESTS); + // each thread has different kernel sizes so different results + const int kernelSizes[] = {3, 5, 7, 9, 11, 13, 15, 17}; + for (int i = 0; i < TOTAL_TESTS; ++i) { + int kernelSize = kernelSizes[i % 8]; + threadTests.push_back(std::make_unique(batchSize, width, height, format, device, kernelSize, i)); + } + + AverageBlur op(17, 17); // shared op for all threads + std::vector threads; + std::atomic nextTestIndex{0}; + + auto threadFunc = [&]() { + while (true) { + int idx = nextTestIndex.fetch_add(1); + if (idx >= TOTAL_TESTS) break; + ThreadTest* test = threadTests[idx].get(); + // All threads call the same operator instance concurrently + op(test->stream, test->input, test->output, test->ksize, test->ksize, -1, -1, BorderMode, device); + } + }; + + for (int i = 0; i < NUM_THREADS; i++) { + threads.emplace_back(threadFunc); + } + for (auto& thread : threads) { + thread.join(); + } + for (auto& threadTest : threadTests) { + HIP_VALIDATE_NO_ERRORS(hipStreamSynchronize(threadTest->stream)); + } + + for (auto& threadTest : threadTests) { + std::vector outputData(threadTest->output.shape().size()); + CopyTensorIntoVector(outputData, threadTest->output); + + int anchor = threadTest->ksize >> 1; + std::vector ref = GenerateGoldenAverageBlur( + threadTest->inputData, batchSize, width, height, threadTest->ksize, threadTest->ksize, anchor, anchor); + + CompareVectorsNear(outputData, ref, 1); + } +} + +void TestNegativeAverageBlur() { + TensorShape validShape(TensorLayout(eTensorLayout::TENSOR_LAYOUT_NHWC), {1, 1, 1, 1}); + Tensor validGPUTensor(validShape, DataType(eDataType::DATA_TYPE_U8), eDeviceType::GPU); + Tensor validCPUTensor(validShape, DataType(eDataType::DATA_TYPE_U8), eDeviceType::CPU); + + AverageBlur op(3, 3); + + { + // Test wrong device + EXPECT_EXCEPTION( + op(nullptr, validCPUTensor, validGPUTensor, 3, 3, 1, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_OPERATION); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validCPUTensor, 3, 3, 1, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); + } + + { + // Test unsupported/mismatch input/output data type + Tensor invalidTensor(validGPUTensor.shape(), DataType(eDataType::DATA_TYPE_U32), eDeviceType::GPU); + EXPECT_EXCEPTION( + op(nullptr, invalidTensor, validGPUTensor, 3, 3, -1, -1, BORDER_TYPE_REFLECT, eDeviceType::GPU), + eStatusType::NOT_IMPLEMENTED); + EXPECT_EXCEPTION( + op(nullptr, validCPUTensor, invalidTensor, 3, 3, -1, -1, BORDER_TYPE_REFLECT, eDeviceType::CPU), + eStatusType::INVALID_COMBINATION); + Tensor validGPUS16Tensor(validShape, DataType(eDataType::DATA_TYPE_S16), eDeviceType::GPU); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUS16Tensor, 3, 3, -1, -1, BORDER_TYPE_REFLECT, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); + } + + { + // Test unsupported input/output layout + TensorShape invalidLayoutShape(TensorLayout(eTensorLayout::TENSOR_LAYOUT_NC), {1, 1}); + Tensor invalidTensor(invalidLayoutShape, DataType(eDataType::DATA_TYPE_U8), eDeviceType::GPU); + EXPECT_EXCEPTION(op(nullptr, invalidTensor, validGPUTensor, 3, 3, -1, -1, BORDER_TYPE_WRAP, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, invalidTensor, 3, 3, -1, -1, BORDER_TYPE_WRAP, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); + } + + { + // Test input/output shape mismatch + Tensor invalidTensor(TensorShape(validGPUTensor.layout(), {2, 2, 2, 2}), DataType(eDataType::DATA_TYPE_U8), + eDeviceType::GPU); + EXPECT_EXCEPTION( + op(nullptr, invalidTensor, validGPUTensor, 3, 3, -1, -1, BORDER_TYPE_REPLICATE, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); + } + + { + // Test bad op construction (bad max ksize) + EXPECT_EXCEPTION(AverageBlur opBadW(0, 3), eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(AverageBlur opBadH(3, -3), eStatusType::INVALID_VALUE); + } + + { + // Test bad kernel size + // exceeding max + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 3, 5, -1, -1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 5, 3, -1, -1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + // not odd + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 1, 2, -1, -1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 2, 1, -1, -1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + // not positive + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 0, 2, -1, -1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 2, -1, -1, -1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + } + + { + // Test bad anchor + // >= ksize + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 3, 3, 4, -1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 3, 3, -1, 3, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + // Negative and not -1 + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 3, 3, 2, -2, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 3, 3, -2, 2, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + } +} +} // namespace + +int main(int argc, char** argv) { + (void)argc; + (void)argv; + TEST_CASES_BEGIN(); + + // Test negative operator cases + TEST_CASE(TestNegativeAverageBlur()); + + // Test concurrency on GPU and CPU + TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_F32, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_F32, eDeviceType::CPU))); + + // GPU correctness tests + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U8, 3, 3, -1, -1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U16, 3, 3, 0, 0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_U16, 5, 5, 2, 2, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_S16, 3, 3, 1, 1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 32, 32, FMT_S32, 5, 3, 4, 0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 32, 32, FMT_S32, 3, 3, -1, -1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 7, 7, 3, 6, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_F32, 5, 5, 0, 0, eDeviceType::GPU))); + + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGB8, 5, 5, -1, -1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB8, 3, 3, 2, 2, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB16, 3, 3, 0, 1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGBs16, 7, 3, 6, 1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGBs16, 3, 3, 1, 0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_RGBs32, 5, 5, 2, 4, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBf32, 3, 3, -1, -1, eDeviceType::GPU))); + + TEST_CASE((TestCorrectness(1, 10, 10, FMT_RGBA8, 3, 3, 0, 2, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(5, 64, 64, FMT_RGBA8, 7, 7, 3, 3, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGBA16, 3, 3, -1, -1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGBA16, 5, 3, 2, 0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGBAs16, 3, 3, 1, 1, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_RGBAs32, 5, 5, 0, 4, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBAf32, 5, 3, 2, 1, eDeviceType::GPU))); + + // CPU correctness tests + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U8, 3, 3, -1, -1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U16, 3, 3, 0, 0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_U16, 5, 5, 2, 2, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_S16, 3, 3, 1, 1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 32, 32, FMT_S32, 5, 3, 4, 0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 32, 32, FMT_S32, 3, 3, -1, -1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 7, 7, 3, 6, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_F32, 5, 5, 0, 0, eDeviceType::CPU))); + + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGB8, 5, 5, -1, -1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB8, 3, 3, 2, 2, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB16, 3, 3, 0, 1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGBs16, 7, 3, 6, 1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGBs16, 3, 3, 1, 0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_RGBs32, 5, 5, 2, 4, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBf32, 3, 3, -1, -1, eDeviceType::CPU))); + + TEST_CASE((TestCorrectness(1, 10, 10, FMT_RGBA8, 3, 3, 0, 2, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(5, 64, 64, FMT_RGBA8, 7, 7, 3, 3, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGBA16, 3, 3, -1, -1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGBA16, 5, 3, 2, 0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGBAs16, 3, 3, 1, 1, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_RGBAs32, 5, 5, 0, 4, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBAf32, 5, 3, 2, 1, eDeviceType::CPU))); + + TEST_CASES_END(); +} diff --git a/tests/roccv/python/test_op_average_blur.py b/tests/roccv/python/test_op_average_blur.py new file mode 100644 index 00000000..a6ee8ef8 --- /dev/null +++ b/tests/roccv/python/test_op_average_blur.py @@ -0,0 +1,51 @@ +# ############################################################################## +# Copyright (c) - 2026 Advanced Micro Devices, Inc. +# +# 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. +# +# ############################################################################## + +import pytest +import rocpycv + +from test_helpers import compare_tensors, generate_tensor + +@pytest.mark.parametrize("device", [rocpycv.eDeviceType.GPU, rocpycv.eDeviceType.CPU]) +@pytest.mark.parametrize("dtype", [rocpycv.eDataType.U8, rocpycv.eDataType.U16, rocpycv.eDataType.S16, rocpycv.eDataType.S32, rocpycv.eDataType.F32]) +@pytest.mark.parametrize("border_mode", [rocpycv.eBorderType.CONSTANT, rocpycv.eBorderType.WRAP, rocpycv.eBorderType.REFLECT, rocpycv.eBorderType.REFLECT101, rocpycv.eBorderType.REPLICATE]) +@pytest.mark.parametrize("kernel_size, anchor", [ + ((3, 3), (-1, -1)), + ((5, 3), (1, 0)), +]) +@pytest.mark.parametrize("channels", [1, 3, 4]) +@pytest.mark.parametrize("samples,height,width", [ + (1, 56, 24), + (3, 14, 40), + (7, 45, 105) +]) +def test_op_averageblur(samples, height, width, channels, border_mode, kernel_size, anchor, dtype, device): + input = generate_tensor(samples, width, height, channels, dtype, device) + output_golden = rocpycv.Tensor([samples, height, width, channels], rocpycv.eTensorLayout.NHWC, dtype, device) + + stream = rocpycv.Stream() + rocpycv.averageblur_into(output_golden, input, kernel_size, anchor, border_mode, stream=stream, device=device) + output = rocpycv.averageblur(input, kernel_size, anchor, border_mode, stream=stream, device=device) + stream.synchronize() + + compare_tensors(output, output_golden) \ No newline at end of file