diff --git a/CHANGELOG.md b/CHANGELOG.md index 245c5fca..5e89b01c 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 `Laplacian` operator. - Dockerfile support for running rocCV in a container environment. ### Changed diff --git a/benchmarks/src/roccv/bench_laplacian.cpp b/benchmarks/src/roccv/bench_laplacian.cpp new file mode 100644 index 00000000..8fb3a5f0 --- /dev/null +++ b/benchmarks/src/roccv/bench_laplacian.cpp @@ -0,0 +1,101 @@ +/* + * 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 RunLaplacianBenchmark(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 ksize = roccvbench::GetParamValue(params, "ksize"); + float scale = roccvbench::GetParamValue(params, "scale"); + 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); + + Laplacian op; + + hipStream_t stream; + HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream)); + + roccvbench::RecordRuns(stream, runs, warmupRuns, results.executionTimes, + [&]() { op(stream, input, output, ksize, scale, borderType, DeviceType); }); + + HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream)); + + return results; +} + +#define DEFINE_LAPLACIAN_BENCHMARK(name, device, inFormat, outFormat, ksize, scale, borderType) \ + BENCHMARK_P(Laplacian, name, \ + BENCH_PARAMS(BENCH_PARAM("inFormat", inFormat), BENCH_PARAM("outFormat", outFormat), \ + BENCH_PARAM("ksize", ksize), BENCH_PARAM("scale", scale), \ + BENCH_PARAM("borderType", borderType))) { \ + return RunLaplacianBenchmark(params); \ + } + +// GPU benchmarks +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 1, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 3, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 1, -1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, 3, -1.0f, eBorderType::BORDER_TYPE_REFLECT); + +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 1, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 3, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 1, -1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, 3, -1.0f, eBorderType::BORDER_TYPE_REFLECT); + +// CPU benchmarks +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 1, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 3, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 1, -1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, 3, -1.0f, eBorderType::BORDER_TYPE_REFLECT); + +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 1, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 3, 1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 1, -1.0f, eBorderType::BORDER_TYPE_REFLECT); +DEFINE_LAPLACIAN_BENCHMARK(CPU, eDeviceType::CPU, FMT_U8, FMT_U8, 3, -1.0f, 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..c0c7b0a1 100644 --- a/docs/reference/rocCV-supported-operators.rst +++ b/docs/reference/rocCV-supported-operators.rst @@ -23,6 +23,7 @@ The rocCV is a collection of the following computer vision operators that are su "GammaContrast","Adjusts the gamma contrast on images in a tensor.","U8","NHWC, HWC" "Gaussian","Applies a Gaussian filter.","U8, U16, S16, S32, F32", "NHWC, HWC" "Histogram","Calculates a histogram of values from a grayscale image.","U8","NHWC, HWC" + "Laplacian", "Applies a Laplacian filter","U8, U16, F32", "NHWC, HWC" "NonMaximumSuppression","Performs non-maximum suppression on batches of bounding boxes based on a score and IoU threshold.","S16, 4S16","NW, NWC" "Normalize","Normalizes an input tensor using a provided mean and standard deviation.","U8, S8, F32","NHWC, HWC" "Remap","Maps pixels in an image from one projection to another projection in a new image.","U8","NHWC, HWC" diff --git a/docs/supported-operators.md b/docs/supported-operators.md index 78ba9b86..cd8a1da4 100644 --- a/docs/supported-operators.md +++ b/docs/supported-operators.md @@ -18,6 +18,7 @@ See below for a list of Computer Vision operators rocCV supports. |GammaContrast|Adjusts the gamma contrast of an image.|U8, U16, U32, F32|NHWC, HWC|Both| |Gaussian|Applies a Gaussian filter for image smoothing.|U8, U16, S16, S32, F32|NHWC, HWC|Both| |Histogram|Calculates a histogram of values from a grayscale image.|U8|NHWC, HWC|Both| +|Laplacian|Applies a Laplacian filter.|U8, U16, F32|NHWC, HWC|Both| |NonMaximumSuppression|Performs non-maximum suppression on batches of bounding boxes based on a score and IoU threshold.|S16, 4S16|NW, NWC|Both| |Normalize|Normalizes image pixels' range using the provided shift and scale parameters.|U8, S8, U16, S16, U32, S32, F32|NHWC, HWC|Both| |Reformat|Converts a tensor between different memory layouts (e.g., NHWC, NCHW, HWC, CHW).|U8, S8, U16, S16, U32, S32, F32, F64|NHWC, NCHW, HWC, CHW|Both| diff --git a/include/op_laplacian.hpp b/include/op_laplacian.hpp new file mode 100644 index 00000000..bc46f6d0 --- /dev/null +++ b/include/op_laplacian.hpp @@ -0,0 +1,78 @@ +/** +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 + +#include "core/tensor.hpp" + +namespace roccv { +/** + * @brief Class for managing the Laplacian operator. + * + */ +class Laplacian final : public IOperator { + public: + /** + * @brief Construct a new Laplacian object. + * + * Limitations: + * + * Input: + * Supported TensorLayout(s): [HWC, NHWC] + * Channels: [1, 3, 4] + * Supported DataType(s): [U8, U16, F32] + * + * Output: + * Supported TensorLayout(s): [HWC, NHWC] + * Channels: [1, 3, 4] + * Supported DataType(s): [U8, U16, F32] + * + * Parameter requirements: + * ksize: Must be 1 or 3. + * + * 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 operator on. + * @param[in] input Input tensor with image data. + * @param[out] output Output tensor for storing modified image data. + * @param[in] ksize Aperture size used to compute the second derivative filters. Must be 1 or 3. + * @param[in] scale Scale factor for the Laplacian values. + * @param[in] borderMode A border type to identify the pixel extrapolation method. + * @param[in] device The device to run this operator on. (Default: GPU) + */ + void operator()(hipStream_t stream, const roccv::Tensor &input, const roccv::Tensor &output, int32_t ksize, + float scale, eBorderType borderMode, eDeviceType device = eDeviceType::GPU) const; +}; +} // namespace roccv \ No newline at end of file diff --git a/include/roccv_operators.hpp b/include/roccv_operators.hpp index d3791de7..e2146db6 100644 --- a/include/roccv_operators.hpp +++ b/include/roccv_operators.hpp @@ -34,6 +34,7 @@ THE SOFTWARE. #include "op_gamma_contrast.hpp" #include "op_gaussian.hpp" #include "op_histogram.hpp" +#include "op_laplacian.hpp" #include "op_non_max_suppression.hpp" #include "op_normalize.hpp" #include "op_remap.hpp" diff --git a/python/include/operators/py_op_laplacian.hpp b/python/include/operators/py_op_laplacian.hpp new file mode 100644 index 00000000..63f9ac8c --- /dev/null +++ b/python/include/operators/py_op_laplacian.hpp @@ -0,0 +1,40 @@ +/** +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 PyOpLaplacian { + public: + static void Export(py::module& m); + static PyTensor Execute(PyTensor& input, int ksize, float scale, eBorderType borderMode, + std::optional> stream, eDeviceType device); + static void ExecuteInto(PyTensor& output, PyTensor& input, int ksize, float scale, 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..00bec975 100644 --- a/python/src/main.cpp +++ b/python/src/main.cpp @@ -37,6 +37,7 @@ THE SOFTWARE. #include "operators/py_op_gamma_contrast.hpp" #include "operators/py_op_gaussian.hpp" #include "operators/py_op_histogram.hpp" +#include "operators/py_op_laplacian.hpp" #include "operators/py_op_non_max_suppression.hpp" #include "operators/py_op_normalize.hpp" #include "operators/py_op_reformat.hpp" @@ -80,6 +81,7 @@ PYBIND11_MODULE(rocpycv, m) { PyOpBrightnessContrast::Export(m); PyOpGammaContrast::Export(m); PyOpGaussian::Export(m); + PyOpLaplacian::Export(m); PyOpComposite::Export(m); PyOpCopyMakeBorder::Export(m); PyOpCenterCrop::Export(m); diff --git a/python/src/operators/py_op_laplacian.cpp b/python/src/operators/py_op_laplacian.cpp new file mode 100644 index 00000000..beddcfca --- /dev/null +++ b/python/src/operators/py_op_laplacian.cpp @@ -0,0 +1,93 @@ +/** +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_laplacian.hpp" + +#include + +#include "py_helpers.hpp" + +PyTensor PyOpLaplacian::Execute(PyTensor& input, int ksize, float scale, 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); + + roccv::Laplacian op; + op(hipStream, *inputTensor, *outputTensor, ksize, scale, borderMode, device); + return PyTensor(outputTensor); +} + +void PyOpLaplacian::ExecuteInto(PyTensor& output, PyTensor& input, int ksize, float scale, eBorderType borderMode, + std::optional> stream, eDeviceType device) { + hipStream_t hipStream = stream.has_value() ? stream.value().get().getStream() : nullptr; + + roccv::Laplacian op; + op(hipStream, *input.getTensor(), *output.getTensor(), ksize, scale, borderMode, device); +} + +void PyOpLaplacian::Export(py::module& m) { + using namespace py::literals; + m.def("laplacian", &PyOpLaplacian::Execute, "src"_a, "ksize"_a, "scale"_a = 1.0f, + "borderMode"_a = BORDER_TYPE_CONSTANT, py::kw_only(), "stream"_a = nullptr, "device"_a = eDeviceType::GPU, + R"pbdoc( + + Executes the Laplacian 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. + ksize (int): Aperture size used to compute the second-derivative filters. Must be 1 or 3. + scale (float, optional): Scale factor for the Laplacian values. Defaults to 1 (no scale). + 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("laplacian_into", &PyOpLaplacian::ExecuteInto, "dst"_a, "src"_a, "ksize"_a, "scale"_a = 1.0f, + "borderMode"_a = BORDER_TYPE_CONSTANT, py::kw_only(), "stream"_a = nullptr, "device"_a = eDeviceType::GPU, + R"pbdoc( + + Executes the Laplacian 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. + ksize (int): Aperture size used to compute the second-derivative filters. Must be 1 or 3. + scale (float, optional): Scale factor for the Laplacian values. Defaults to 1 (no scale). + 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_laplacian.cpp b/src/op_laplacian.cpp new file mode 100644 index 00000000..1c61639f --- /dev/null +++ b/src/op_laplacian.cpp @@ -0,0 +1,115 @@ +/** +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_laplacian.hpp" + +#include + +#include + +#include "common/validation_helpers.hpp" +#include "core/detail/casting.hpp" +#include "details/filter_2d.hpp" + +namespace { +using namespace roccv; +constexpr int LaplaceKWidth = 3; +constexpr int LaplaceKHeight = 3; + +class LaplacianKernel { + public: + LaplacianKernel(const float (&kernel)[9]) { + for (int i = 0; i < 9; i++) { + m_kernel[i] = kernel[i]; + } + } + LaplacianKernel& operator*=(float scale) { + for (int i = 0; i < 9; i++) { + m_kernel[i] *= scale; + } + return *this; + } + __device__ __host__ const float& operator[](int i) const { return m_kernel[i]; } + + private: + float m_kernel[9] = {}; +}; +// clang-format off +const LaplacianKernel LK1{ + {0.0f, 1.0f, 0.0f, + 1.0f, -4.0f, 1.0f, + 0.0f, 1.0f, 0.0f} +}; +const LaplacianKernel LK3{ + {2.0f, 0.0f, 2.0f, + 0.0f, -8.0f, 0.0f, + 2.0f, 0.0f, 2.0f} +}; +// clang-format on +} // namespace + +namespace roccv { +void Laplacian::operator()(hipStream_t stream, const roccv::Tensor& input, const roccv::Tensor& output, int32_t ksize, + float scale, eBorderType borderMode, eDeviceType device) const { + // Validate input tensor + CHECK_TENSOR_DEVICE(input, device); + CHECK_TENSOR_DATATYPES(input, DATA_TYPE_U8, DATA_TYPE_U16, 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 ksize + if (!(ksize == 1 || ksize == 3)) { + throw roccv::Exception("Invalid ksize = " + std::to_string(ksize) + ": Must be 1 or 3.", + eStatusType::INVALID_VALUE); + } + + // get the right kernel + LaplacianKernel kernel = (ksize == 1) ? LK1 : LK3; + if (scale != 1) { + kernel *= scale; + } + + // compute the anchor to be center of kernel + int anchorX = -1; + int anchorY = -1; + processAnchor(anchorX, anchorY, LaplaceKWidth, LaplaceKHeight); + + // 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_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); + + func(stream, input, output, kernel, LaplaceKWidth, LaplaceKHeight, anchorX, anchorY, borderMode, device); +} +} // namespace roccv \ No newline at end of file diff --git a/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp new file mode 100644 index 00000000..cb86feb5 --- /dev/null +++ b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp @@ -0,0 +1,235 @@ +/** +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 Laplacian 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] ksize Aperture size for the second derivative filters. Must be 1 or 3. + * @param[in] scale Scale factor for the Laplacian values. + * @return Vector containing the results of the operation. + */ +template > +std::vector GenerateGoldenLaplacian(std::vector& input, int32_t batchSize, int32_t width, int32_t height, + int ksize, float scale) { + 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>; + + std::array kernel; + if (ksize == 1) { + kernel = {0.0f, 1.0f, 0.0f, 1.0f, -4.0f, 1.0f, 0.0f, 1.0f, 0.0f}; + } else if (ksize == 3) { + kernel = {2.0f, 0.0f, 2.0f, 0.0f, -8.0f, 0.0f, 2.0f, 0.0f, 2.0f}; + } + + worktype res; + int kIdx; + for (int b = 0; b < dst.batches(); b++) { + for (int j = 0; j < dst.height(); j++) { + for (int i = 0; i < dst.width(); i++) { + res = SetAll(0); + kIdx = 0; + for (int y = j - 1; y <= j + 1; y++) { + for (int x = i - 1; x <= i + 1; x++) { + res += StaticCast(src.at(b, y, x, 0)) * kernel[kIdx++]; + } + } + dst.at(b, j, i, 0) = SaturateCast(res * scale); + } + } + } + return output; +} + +/** + * @brief Tests correctness of the Laplacian 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] ksize Aperture size for the second derivative filters. Must be 1 or 3. + * @param[in] scale Scale factor for the Laplacian values. + * @param[in] device Device this correctness test should be run on. + */ +template > +void TestCorrectness(int batchSize, int width, int height, ImageFormat format, int ksize, float scale, + 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); + + hipStream_t stream; + HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream)); + Laplacian op; + op(stream, input, output, ksize, scale, 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 = GenerateGoldenLaplacian(inputData, batchSize, width, height, ksize, scale); + + // Compare data in actual output versus the generated golden reference image + CompareVectorsNear(outputData, ref, 1.0); +} + +void TestNegativeLaplacian() { + 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); + + Laplacian op; + + { + // Test wrong device + EXPECT_EXCEPTION(op(nullptr, validCPUTensor, validGPUTensor, 3, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_OPERATION); + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, validCPUTensor, 3, 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, 1, BORDER_TYPE_REFLECT, eDeviceType::GPU), + eStatusType::NOT_IMPLEMENTED); + EXPECT_EXCEPTION(op(nullptr, validCPUTensor, invalidTensor, 3, 1, BORDER_TYPE_REFLECT, eDeviceType::CPU), + eStatusType::INVALID_COMBINATION); + Tensor validGPUU16Tensor(validShape, DataType(eDataType::DATA_TYPE_U16), eDeviceType::GPU); + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, validGPUU16Tensor, 3, 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, 1, BORDER_TYPE_WRAP, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, invalidTensor, 3, 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, 1, BORDER_TYPE_REPLICATE, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); + } + + { + // Test bad ksize + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, validGPUTensor, 4, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, validGPUTensor, 2, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, validGPUTensor, 0, 1, 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(TestNegativeLaplacian()); + + // GPU correctness tests + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U8, 1, 1.0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U16, 3, 1.0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 1, 2.0, eDeviceType::GPU))); + + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGB8, 3, 2.0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB16, 1, -1.0, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBf32, 3, -1.0, eDeviceType::GPU))); + + TEST_CASE((TestCorrectness(5, 64, 64, FMT_RGBA8, 1, 1.5, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGBA16, 3, 1.5, eDeviceType::GPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBAf32, 3, 10, eDeviceType::GPU))); + + // CPU correctness tests + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U8, 1, 1.0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_U16, 3, 1.0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 24, 24, FMT_F32, 1, 2.0, eDeviceType::CPU))); + + TEST_CASE((TestCorrectness(2, 20, 20, FMT_RGB8, 3, 2.0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGB16, 1, -1.0, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBf32, 3, -1.0, eDeviceType::CPU))); + + TEST_CASE((TestCorrectness(5, 64, 64, FMT_RGBA8, 1, 1.5, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(1, 20, 20, FMT_RGBA16, 3, 1.5, eDeviceType::CPU))); + TEST_CASE((TestCorrectness(2, 24, 24, FMT_RGBAf32, 3, 10, eDeviceType::CPU))); + + TEST_CASES_END(); +} \ No newline at end of file diff --git a/tests/roccv/python/test_op_laplacian.py b/tests/roccv/python/test_op_laplacian.py new file mode 100644 index 00000000..75202b27 --- /dev/null +++ b/tests/roccv/python/test_op_laplacian.py @@ -0,0 +1,49 @@ +# ############################################################################## +# 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.F32]) +@pytest.mark.parametrize("border_mode", [rocpycv.eBorderType.CONSTANT, rocpycv.eBorderType.WRAP, rocpycv.eBorderType.REFLECT, rocpycv.eBorderType.REFLECT101, rocpycv.eBorderType.REPLICATE]) +@pytest.mark.parametrize("ksize", [1, 3]) +@pytest.mark.parametrize("scale", [1.0, -1.0, 2.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_laplacian(samples, height, width, channels, border_mode, ksize, scale, 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.laplacian_into(output_golden, input, ksize, scale, border_mode, stream=stream, device=device) + output = rocpycv.laplacian(input, ksize, scale, border_mode, stream=stream, device=device) + stream.synchronize() + + compare_tensors(output, output_golden)