[New Operator] Initial Implementation of Laplacian

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

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 <core/hip_assert.h>
+
+#include <cfenv>
+#include <core/image_format.hpp>
+#include <core/tensor.hpp>
+#include <op_laplacian.hpp>
+#include <roccvbench/registry.hpp>
+#include <roccvbench/utils.hpp>
+
+#include "roccv_bench_helpers.hpp"
+
+using namespace roccv;
+
+template <eDeviceType DeviceType>
+static roccvbench::BenchmarkResults RunLaplacianBenchmark(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 ksize = roccvbench::GetParamValue<int>(params, "ksize");
+    float scale = roccvbench::GetParamValue<float>(params, "scale");
+    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);
+
+    Laplacian op;
+
+    hipStream_t stream;
+    HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream));
+
+    roccvbench::RecordRuns<DeviceType>(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<device>(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);

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"

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|

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 <hip/hip_runtime.h>
+#include <operator_types.h>
+
+#include <i_operator.hpp>
+
+#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

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"

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 <operator_types.h>
+#include <pybind11/pybind11.h>
+
+#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<std::reference_wrapper<PyStream>> stream, eDeviceType device);
+    static void ExecuteInto(PyTensor& output, PyTensor& input, int ksize, float scale, eBorderType borderMode,
+                            std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device);
+};

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);

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 <op_laplacian.hpp>
+
+#include "py_helpers.hpp"
+
+PyTensor PyOpLaplacian::Execute(PyTensor& input, int ksize, float scale, eBorderType borderMode,
+                                std::optional<std::reference_wrapper<PyStream>> stream, eDeviceType device) {
+    hipStream_t hipStream = stream.has_value() ? stream.value().get().getStream() : nullptr;
+
+    auto inputTensor = input.getTensor();
+    auto outputTensor = std::make_shared<roccv::Tensor>(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<std::reference_wrapper<PyStream>> 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");
+}