From 5d19381572d3e7785af2d452afb2bac3648b2945 Mon Sep 17 00:00:00 2001 From: daniellegillai Date: Thu, 16 Jul 2026 10:21:49 -0700 Subject: [PATCH 1/7] filter 2d files --- include/details/filter_2d.hpp | 165 ++++++++++++++++++++ include/kernels/device/filter_2d_device.hpp | 138 ++++++++++++++++ include/kernels/host/filter_2d_host.hpp | 104 ++++++++++++ 3 files changed, 407 insertions(+) create mode 100644 include/details/filter_2d.hpp create mode 100644 include/kernels/device/filter_2d_device.hpp create mode 100644 include/kernels/host/filter_2d_host.hpp diff --git a/include/details/filter_2d.hpp b/include/details/filter_2d.hpp new file mode 100644 index 00000000..c51a7fd0 --- /dev/null +++ b/include/details/filter_2d.hpp @@ -0,0 +1,165 @@ +/** +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/vector_utils.hpp" +#include "core/exception.hpp" +#include "core/tensor.hpp" +#include "core/wrappers/border_wrapper.hpp" +#include "core/wrappers/image_wrapper.hpp" +#include "kernels/device/filter_2d_device.hpp" +#include "kernels/host/filter_2d_host.hpp" +#include "operator_types.h" + +namespace roccv { + +inline void processAnchor(int& anchorX, int& anchorY, int kernelWidth, int kernelHeight) { + if (anchorX < 0) { + anchorX = kernelWidth >> 1; + } + if (anchorY < 0) { + anchorY = kernelHeight >> 1; + } +} + +template +void dispatch_filter_2d_bordertype(hipStream_t stream, const Tensor& input, const Tensor& output, KT kernel, + int kernelWidth, int kernelHeight, int anchorX, int anchorY, eDeviceType device) { + BorderWrapper inputWrapper(input, roccv::detail::SetAll
(0)); + ImageWrapper
outputWrapper(output); + + switch (device) { + case eDeviceType::GPU: { + dim3 block(16, 16); + dim3 grid((outputWrapper.width() + block.x - 1) / block.x, (outputWrapper.height() + block.y - 1) / block.y, + outputWrapper.batches()); + Kernels::Device::filter_2d<<>>(inputWrapper, outputWrapper, kernel, kernelWidth, + kernelHeight, anchorX, anchorY); + break; + } + case eDeviceType::CPU: { + Kernels::Host::filter_2d(inputWrapper, outputWrapper, kernel, kernelWidth, kernelHeight, anchorX, anchorY); + break; + } + } +} +template +void dispatch_filter_2d_bordertype_separable(hipStream_t stream, const Tensor& input, const Tensor& output, KT kernelH, + KT kernelV, int kernelWidth, int kernelHeight, int anchorX, int anchorY, + eDeviceType device) { + BorderWrapper inputWrapper(input, roccv::detail::SetAll
(0)); + ImageWrapper
outputWrapper(output); + Tensor interm(output.shape(), output.dtype(), device); + ImageWrapper
intermWrapper(interm); + + switch (device) { + case eDeviceType::CPU: { + Kernels::Host::filter_2d_horizontal(inputWrapper, intermWrapper, kernelH, kernelWidth, anchorX); + BorderWrapper intermWrapperWithBorder(interm, roccv::detail::SetAll
(0)); + Kernels::Host::filter_2d_vertical(intermWrapperWithBorder, outputWrapper, kernelV, kernelHeight, anchorY); + break; + } + case eDeviceType::GPU: { + // constants copied from Pavel + constexpr int BLOCK_WIDTH = 128; + constexpr int BLOCK_HEIGHT = 128; + + // horizontal pass + { + dim3 block(BLOCK_WIDTH, 1); + dim3 grid((outputWrapper.width() + BLOCK_WIDTH - 1) / BLOCK_WIDTH, outputWrapper.height(), + outputWrapper.batches()); + + int halo = kernelWidth - 1; + int tileWidth = BLOCK_WIDTH + halo; + size_t smemSize = tileWidth * sizeof(DT); + + Kernels::Device::filter_2d_horizontal, ImageWrapper
, KT> + <<>>(inputWrapper, intermWrapper, kernelH, kernelWidth, anchorX); + } + + // Vertical pass + BorderWrapper intermWrapperWithBorder(interm, roccv::detail::SetAll
(0)); + { + dim3 block(1, BLOCK_HEIGHT); + dim3 grid(outputWrapper.width(), (outputWrapper.height() + BLOCK_HEIGHT - 1) / BLOCK_HEIGHT, + outputWrapper.batches()); + + int halo = kernelHeight - 1; + int tileHeight = BLOCK_HEIGHT + halo; + size_t smemSize = tileHeight * sizeof(DT); + + Kernels::Device::filter_2d_vertical, ImageWrapper
, KT> + <<>>(intermWrapperWithBorder, outputWrapper, kernelV, kernelHeight, + anchorY); + } + break; + } + } +} + +template +void dispatch_filter_2d_dtype(hipStream_t stream, const Tensor& input, const Tensor& output, KT kernel, int kernelWidth, + int kernelHeight, int anchorX, int anchorY, eBorderType borderMode, eDeviceType device) { + // clang-format off + static const std::unordered_map> + funcs = { + {eBorderType::BORDER_TYPE_REPLICATE, dispatch_filter_2d_bordertype}, + {eBorderType::BORDER_TYPE_CONSTANT, dispatch_filter_2d_bordertype}, + {eBorderType::BORDER_TYPE_REFLECT, dispatch_filter_2d_bordertype}, + {eBorderType::BORDER_TYPE_REFLECT101, dispatch_filter_2d_bordertype}, + {eBorderType::BORDER_TYPE_WRAP, dispatch_filter_2d_bordertype} + }; + // clang-format on + if (!funcs.contains(borderMode)) { + throw Exception("The given border mode is not supported.", eStatusType::NOT_IMPLEMENTED); + } + + auto func = funcs.at(borderMode); + func(stream, input, output, kernel, kernelWidth, kernelHeight, anchorX, anchorY, device); +} + +template +void dispatch_filter_2d_dtype_separable(hipStream_t stream, const Tensor& input, const Tensor& output, KT kernelH, + KT kernelV, int kernelWidth, int kernelHeight, int anchorX, int anchorY, + eBorderType borderMode, eDeviceType device) { + // clang-format off + static const std::unordered_map> + funcs = { + {eBorderType::BORDER_TYPE_REPLICATE, dispatch_filter_2d_bordertype_separable}, + {eBorderType::BORDER_TYPE_CONSTANT, dispatch_filter_2d_bordertype_separable}, + {eBorderType::BORDER_TYPE_REFLECT, dispatch_filter_2d_bordertype_separable}, + {eBorderType::BORDER_TYPE_REFLECT101, dispatch_filter_2d_bordertype_separable}, + {eBorderType::BORDER_TYPE_WRAP, dispatch_filter_2d_bordertype_separable} + }; + // clang-format on + if (!funcs.contains(borderMode)) { + throw Exception("The given border mode is not supported.", eStatusType::NOT_IMPLEMENTED); + } + + auto func = funcs.at(borderMode); + func(stream, input, output, kernelH, kernelV, kernelWidth, kernelHeight, anchorX, anchorY, device); +} +} // namespace roccv \ No newline at end of file diff --git a/include/kernels/device/filter_2d_device.hpp b/include/kernels/device/filter_2d_device.hpp new file mode 100644 index 00000000..a7ad93ce --- /dev/null +++ b/include/kernels/device/filter_2d_device.hpp @@ -0,0 +1,138 @@ +/** +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 "core/detail/casting.hpp" +#include "core/detail/type_traits.hpp" +#include "core/detail/vector_utils.hpp" +#include "core/wrappers/image_wrapper.hpp" + +namespace Kernels { +namespace Device { +template +__global__ void filter_2d(SrcWrapper input, DstWrapper output, KernelWrapper kernel, int kernelWidth, int kernelHeight, + int anchorX, int anchorY) { + using namespace roccv::detail; + using dst_type = typename DstWrapper::ValueType; + using work_type = MakeType>; + work_type result = SetAll(0); + + const int x = threadIdx.x + blockIdx.x * blockDim.x; + const int y = threadIdx.y + blockIdx.y * blockDim.y; + const int batch = blockIdx.z; + + if (x >= output.width() || y >= output.height()) return; + + // kernel is vector or 1d in row major order + int kernelIndex = 0; + int3 coord{x, y, batch}; + + for (int i = 0; i < kernelHeight; ++i) { + coord.y = y - anchorY + i; + for (int j = 0; j < kernelWidth; ++j) { + coord.x = x - anchorX + j; + result = result + StaticCast(input.at(coord.z, coord.y, coord.x, 0)) * kernel[kernelIndex++]; + } + } + output.at(batch, y, x, 0) = SaturateCast(result); +} + +template +__global__ void filter_2d_horizontal(SrcWrapper input, DstWrapper output, KernelWrapper kernel, int kernelWidth, + int anchorX) { + using namespace roccv::detail; + using work_type = MakeType>; + work_type result = SetAll(0); + + const int x = threadIdx.x + blockIdx.x * blockDim.x; + const int y = blockIdx.y; + const int b = blockIdx.z; + + // smem size is BLOCK_WIDTH + halo on each side + extern __shared__ char smem[]; + T* tile = reinterpret_cast(smem); + + const int halo = kernelWidth - 1; + const int tileWidth = BLOCK_WIDTH + halo; + + // load into shared memory + for (int i = threadIdx.x; i < tileWidth; i += blockDim.x) { + int srcX = blockIdx.x * BLOCK_WIDTH + i - anchorX; + tile[i] = input.at(b, y, srcX, 0); + } + + __syncthreads(); + + if (x >= output.width()) { + return; + } + + int tileIdx = threadIdx.x + anchorX; + for (int kx = 0; kx < kernelWidth; ++kx) { + result = result + StaticCast(tile[tileIdx - anchorX + kx]) * kernel[kx]; + } + + output.at(b, y, x, 0) = SaturateCast(result); +} + +template +__global__ void filter_2d_vertical(SrcWrapper input, DstWrapper output, KernelWrapper kernel, int kernelHeight, + int anchorY) { + using namespace roccv::detail; + using work_type = MakeType>; + work_type result = SetAll(0); + + const int x = blockIdx.x; + const int y = threadIdx.y + blockIdx.y * blockDim.y; + const int b = blockIdx.z; + + // smem size is BLOCK_HEIGHT + halo on each side + extern __shared__ char smem[]; + T* tile = reinterpret_cast(smem); + + const int halo = kernelHeight - 1; + const int tileHeight = BLOCK_HEIGHT + halo; + + // load into shared memory + for (int i = threadIdx.y; i < tileHeight; i += blockDim.y) { + int srcY = blockIdx.y * BLOCK_HEIGHT + i - anchorY; + tile[i] = input.at(b, srcY, x, 0); + } + + __syncthreads(); + + if (y >= output.height()) { + return; + } + + int tileIdx = threadIdx.y + anchorY; + for (int ky = 0; ky < kernelHeight; ++ky) { + result = result + StaticCast(tile[tileIdx - anchorY + ky]) * kernel[ky]; + } + + output.at(b, y, x, 0) = SaturateCast(result); +} +} // namespace Device +} // namespace Kernels \ No newline at end of file diff --git a/include/kernels/host/filter_2d_host.hpp b/include/kernels/host/filter_2d_host.hpp new file mode 100644 index 00000000..e4d78286 --- /dev/null +++ b/include/kernels/host/filter_2d_host.hpp @@ -0,0 +1,104 @@ +/** +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 "core/detail/casting.hpp" +#include "core/detail/type_traits.hpp" +#include "core/detail/vector_utils.hpp" +#include "core/wrappers/image_wrapper.hpp" + +namespace Kernels { +namespace Host { +template +void filter_2d(SrcWrapper input, DstWrapper output, KernelWrapper kernel, int kernelWidth, int kernelHeight, + int anchorX, int anchorY) { + using namespace roccv::detail; + using dst_type = typename DstWrapper::ValueType; + using work_type = MakeType>; + +#pragma omp parallel for + for (int batch = 0; batch < output.batches(); batch++) { + for (int y = 0; y < output.height(); y++) { + for (int x = 0; x < output.width(); x++) { + work_type result = SetAll(0); + int kernelIndex = 0; + int3 coord{x, y, batch}; + for (int i = 0; i < kernelHeight; ++i) { + coord.y = y - anchorY + i; + for (int j = 0; j < kernelWidth; ++j) { + coord.x = x - anchorX + j; + result = result + + StaticCast(input.at(coord.z, coord.y, coord.x, 0)) * kernel[kernelIndex++]; + } + } + output.at(batch, y, x, 0) = SaturateCast(result); + } + } + } +} + +template +void filter_2d_horizontal(SrcWrapper input, DstWrapper output, KernelWrapper kernel, int kernelWidth, int anchorX) { + using namespace roccv::detail; + using dst_type = typename DstWrapper::ValueType; + using work_type = MakeType>; + +#pragma omp parallel for + for (int batch = 0; batch < output.batches(); batch++) { + for (int y = 0; y < output.height(); ++y) { + for (int x = 0; x < output.width(); ++x) { + work_type result = SetAll(0); + for (int kx = 0; kx < kernelWidth; ++kx) { + int srcX = x - anchorX + kx; + result = result + StaticCast(input.at(batch, y, srcX, 0)) * kernel[kx]; + } + output.at(batch, y, x, 0) = SaturateCast(result); + } + } + } +} + +template +void filter_2d_vertical(SrcWrapper input, DstWrapper output, KernelWrapper kernel, int kernelHeight, int anchorY) { + using namespace roccv::detail; + using dst_type = typename DstWrapper::ValueType; + using work_type = MakeType>; + +#pragma omp parallel for + for (int batch = 0; batch < output.batches(); batch++) { + for (int y = 0; y < output.height(); ++y) { + for (int x = 0; x < output.width(); ++x) { + work_type result = SetAll(0); + for (int ky = 0; ky < kernelHeight; ++ky) { + int srcY = y - anchorY + ky; + result = result + StaticCast(input.at(batch, srcY, x, 0)) * kernel[ky]; + } + output.at(batch, y, x, 0) = SaturateCast(result); + } + } + } +} +} // namespace Host +} // namespace Kernels \ No newline at end of file From 4c00f603c223194b7c5a83a4d3e74d97e775a3c4 Mon Sep 17 00:00:00 2001 From: daniellegillai Date: Thu, 16 Jul 2026 15:33:28 -0700 Subject: [PATCH 2/7] basic implementation and tests passing --- include/kernels/common/laplacian_kernels.hpp | 63 ++++ include/op_laplacian.hpp | 78 ++++ include/roccv_operators.hpp | 1 + src/op_laplacian.cpp | 88 +++++ .../src/tests/operators/test_op_laplacian.cpp | 351 ++++++++++++++++++ 5 files changed, 581 insertions(+) create mode 100644 include/kernels/common/laplacian_kernels.hpp create mode 100644 include/op_laplacian.hpp create mode 100644 src/op_laplacian.cpp create mode 100644 tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp diff --git a/include/kernels/common/laplacian_kernels.hpp b/include/kernels/common/laplacian_kernels.hpp new file mode 100644 index 00000000..2953bcc6 --- /dev/null +++ b/include/kernels/common/laplacian_kernels.hpp @@ -0,0 +1,63 @@ +/** +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 + +namespace Kernels { + +constexpr int LaplaceKWidth = 3; +constexpr int LaplaceKHeight = 3; + +struct LaplacianKernel { + constexpr LaplacianKernel& operator*=(float scale) { + for (int i = 0; i < 9; i++) { + m_kernel[i] *= scale; + } + return *this; + } + + constexpr float& operator[](int i) { + return m_kernel[i]; + } + + constexpr const float& operator[](int i) const { + return m_kernel[i]; + } + + float m_kernel[9] = {}; +}; + +// clang-format off +constexpr LaplacianKernel LK1 { + {0.0f, 1.0f, 0.0f, + 1.0f, -4.0f, 1.0f, + 0.0f, 1.0f, 0.0f} +}; + +constexpr LaplacianKernel LK3 { + {2.0f, 0.0f, 2.0f, + 0.0f, -8.0f, 0.0f, + 2.0f, 0.0f, 2.0f} +}; +// clang-format on + +} // namespace Kernels diff --git a/include/op_laplacian.hpp b/include/op_laplacian.hpp new file mode 100644 index 00000000..48df378e --- /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 Brightness Contrast 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 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 41c47571..44c8dfac 100644 --- a/include/roccv_operators.hpp +++ b/include/roccv_operators.hpp @@ -33,6 +33,7 @@ THE SOFTWARE. #include "op_flip.hpp" #include "op_gamma_contrast.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/src/op_laplacian.cpp b/src/op_laplacian.cpp new file mode 100644 index 00000000..c5df9cad --- /dev/null +++ b/src/op_laplacian.cpp @@ -0,0 +1,88 @@ +/** +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" +#include "kernels/common/laplacian_kernels.hpp" + +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); + } + + using namespace Kernels; + LaplacianKernel kernel; + + if (ksize == 1) { + kernel = LK1; + } else if (ksize == 3) { + kernel = 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_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); + + 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..f05c78a6 --- /dev/null +++ b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp @@ -0,0 +1,351 @@ +/** +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. 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. 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 + // TODO check on delta + CompareVectorsNear(outputData, ref, 1); +} + +/** + * @brief Tests correctness of the Gaussian 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] kernelWidth Kernel width. + * @param[in] kernelHeight Kernel height. + * @param[in] device Device this correctness test should be run on. + */ +/* +template > +void TestCorrectnessConcurrent(int batchSize, int width, int height, ImageFormat format, int kernelWidth, + int kernelHeight, eDeviceType device) { + constexpr int NUM_THREADS = 8; + constexpr int TOTAL_TESTS = 80; + + struct ThreadTest { + Tensor input; + Tensor output; + std::vector inputData; + hipStream_t stream; + double sigma; + + ThreadTest(int b, int w, int h, ImageFormat fmt, eDeviceType dev, double s, int id) + : input(b, {w, h}, fmt, dev), output(b, {w, h}, fmt, dev), inputData(input.shape().size()), sigma(s) { + 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 sigma so different results + for (int i = 0; i < TOTAL_TESTS; ++i) { + double sigma = 0.25 + i * 0.02; + threadTests.push_back(std::make_unique(batchSize, width, height, format, device, sigma, i)); + } + + Gaussian op(kernelWidth, kernelHeight); // 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, kernelWidth, kernelHeight, test->sigma, test->sigma, 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); + + std::vector ref = + GenerateGoldenGaussian(threadTest->inputData, batchSize, width, height, kernelWidth, + kernelHeight, threadTest->sigma, threadTest->sigma); + + CompareVectorsNear(outputData, ref, 1); + } +} + +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); + + Gaussian 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(Gaussian opBadW(0, 3), eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(Gaussian opBadH(3, -3), eStatusType::INVALID_VALUE); + } + + { + // Test bad kernel size (exceeding max) and sigma X + 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); + // Kernel size 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); + // Sigma x not positive + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 3, 3, 0, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), + eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION( + op(nullptr, validGPUTensor, validGPUTensor, 3, 3, -0.1, 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(TestNegativeGaussian()); + + // Test concurrency on GPU and CPU + TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_U8, 7, 7, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_U8, 7, 7, eDeviceType::CPU))); + + TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_F32, 7, 7, eDeviceType::GPU))); + TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_F32, 7, 7, eDeviceType::CPU))); + */ + // 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 From 463386d48c442fcb13e4332bf68ab1ae63ffe481 Mon Sep 17 00:00:00 2001 From: daniellegillai Date: Thu, 16 Jul 2026 16:16:02 -0700 Subject: [PATCH 3/7] removed unused dtypes and minor changes --- src/op_laplacian.cpp | 8 +- .../src/tests/operators/test_op_laplacian.cpp | 169 +++--------------- 2 files changed, 30 insertions(+), 147 deletions(-) diff --git a/src/op_laplacian.cpp b/src/op_laplacian.cpp index c5df9cad..9f58ce15 100644 --- a/src/op_laplacian.cpp +++ b/src/op_laplacian.cpp @@ -22,6 +22,7 @@ THE SOFTWARE. #include "op_laplacian.hpp" #include + #include #include "common/validation_helpers.hpp" @@ -31,7 +32,7 @@ THE SOFTWARE. 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 { + 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); @@ -74,15 +75,12 @@ void Laplacian::operator()(hipStream_t stream, const roccv::Tensor &input, const { {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); - func(stream, input, output, kernel, LaplaceKWidth, LaplaceKHeight, anchorX, - anchorY, borderMode, device); + 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 index f05c78a6..9a40139f 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp @@ -55,7 +55,7 @@ namespace { */ template > std::vector GenerateGoldenLaplacian(std::vector& input, int32_t batchSize, int32_t width, int32_t height, - int ksize, float scale) { + int ksize, float scale) { std::vector output(input.size()); BorderWrapper src(ImageWrapper(input, batchSize, width, height), SetAll(0)); ImageWrapper dst(output, batchSize, width, height); @@ -65,9 +65,9 @@ std::vector GenerateGoldenLaplacian(std::vector& input, int32_t batchSiz 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}; + 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}; + kernel = {2.0f, 0.0f, 2.0f, 0.0f, -8.0f, 0.0f, 2.0f, 0.0f, 2.0f}; } worktype res; @@ -104,8 +104,8 @@ std::vector GenerateGoldenLaplacian(std::vector& input, int32_t batchSiz * @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) { +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); @@ -141,124 +141,40 @@ void TestCorrectness(int batchSize, int width, int height, ImageFormat format, i CompareVectorsNear(outputData, ref, 1); } -/** - * @brief Tests correctness of the Gaussian 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] kernelWidth Kernel width. - * @param[in] kernelHeight Kernel height. - * @param[in] device Device this correctness test should be run on. - */ -/* -template > -void TestCorrectnessConcurrent(int batchSize, int width, int height, ImageFormat format, int kernelWidth, - int kernelHeight, eDeviceType device) { - constexpr int NUM_THREADS = 8; - constexpr int TOTAL_TESTS = 80; - - struct ThreadTest { - Tensor input; - Tensor output; - std::vector inputData; - hipStream_t stream; - double sigma; - - ThreadTest(int b, int w, int h, ImageFormat fmt, eDeviceType dev, double s, int id) - : input(b, {w, h}, fmt, dev), output(b, {w, h}, fmt, dev), inputData(input.shape().size()), sigma(s) { - 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 sigma so different results - for (int i = 0; i < TOTAL_TESTS; ++i) { - double sigma = 0.25 + i * 0.02; - threadTests.push_back(std::make_unique(batchSize, width, height, format, device, sigma, i)); - } - - Gaussian op(kernelWidth, kernelHeight); // 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, kernelWidth, kernelHeight, test->sigma, test->sigma, 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); - - std::vector ref = - GenerateGoldenGaussian(threadTest->inputData, batchSize, width, height, kernelWidth, - kernelHeight, threadTest->sigma, threadTest->sigma); - - CompareVectorsNear(outputData, ref, 1); - } -} - 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); - Gaussian op(3, 3); + Laplacian op; { // 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); + 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, 3, 1, 1, BORDER_TYPE_REFLECT, 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, 3, 1, 1, BORDER_TYPE_REFLECT, eDeviceType::CPU), + 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); - 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), + EXPECT_EXCEPTION(op(nullptr, invalidTensor, validGPUTensor, 3, 1, BORDER_TYPE_WRAP, eDeviceType::GPU), eStatusType::INVALID_COMBINATION); - EXPECT_EXCEPTION(op(nullptr, validGPUTensor, invalidTensor, 3, 3, 1, 1, BORDER_TYPE_WRAP, eDeviceType::GPU), + EXPECT_EXCEPTION(op(nullptr, validGPUTensor, invalidTensor, 3, 1, BORDER_TYPE_WRAP, eDeviceType::GPU), eStatusType::INVALID_COMBINATION); } @@ -266,43 +182,20 @@ void TestNegativeLaplacian() { // 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(Gaussian opBadW(0, 3), eStatusType::INVALID_VALUE); - EXPECT_EXCEPTION(Gaussian opBadH(3, -3), eStatusType::INVALID_VALUE); + EXPECT_EXCEPTION(op(nullptr, invalidTensor, validGPUTensor, 3, 1, BORDER_TYPE_REPLICATE, eDeviceType::GPU), + eStatusType::INVALID_COMBINATION); } { - // Test bad kernel size (exceeding max) and sigma X - 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); - // Kernel size 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); - // Sigma x not positive - EXPECT_EXCEPTION( - op(nullptr, validGPUTensor, validGPUTensor, 3, 3, 0, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), - eStatusType::INVALID_VALUE); - EXPECT_EXCEPTION( - op(nullptr, validGPUTensor, validGPUTensor, 3, 3, -0.1, 1, BORDER_TYPE_CONSTANT, eDeviceType::GPU), - eStatusType::INVALID_VALUE); + // 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) { @@ -310,17 +203,9 @@ int main(int argc, char** argv) { (void)argv; TEST_CASES_BEGIN(); - /* // Test negative operator cases - TEST_CASE(TestNegativeGaussian()); - - // Test concurrency on GPU and CPU - TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_U8, 7, 7, eDeviceType::GPU))); - TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_U8, 7, 7, eDeviceType::CPU))); + TEST_CASE(TestNegativeLaplacian()); - TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_F32, 7, 7, eDeviceType::GPU))); - TEST_CASE((TestCorrectnessConcurrent(1, 64, 64, FMT_F32, 7, 7, eDeviceType::CPU))); - */ // 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))); From 4c65cebd2ecd0207c8f527caa339b589f737e837 Mon Sep 17 00:00:00 2001 From: daniellegillai Date: Fri, 17 Jul 2026 10:08:16 -0700 Subject: [PATCH 4/7] laplacian bench --- benchmarks/src/roccv/bench_laplacian.cpp | 101 +++++++++++++++++++++++ 1 file changed, 101 insertions(+) create mode 100644 benchmarks/src/roccv/bench_laplacian.cpp 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 From 34827c1c43293a5f2e9024ee6c252ca88fc6e132 Mon Sep 17 00:00:00 2001 From: daniellegillai Date: Mon, 20 Jul 2026 10:42:38 -0700 Subject: [PATCH 5/7] changelog and docs --- CHANGELOG.md | 4 +++- docs/reference/rocCV-supported-operators.rst | 1 + docs/supported-operators.md | 1 + 3 files changed, 5 insertions(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 5fe1a756..5e89b01c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -6,7 +6,9 @@ The full documentation for rocCV is available at [https://rocm.docs.amd.com/proj ### Added - Added `Reformat` operator. -- Added `BrightnessContrast` operator +- Added `BrightnessContrast` operator. +- Added `Gaussian` operator. +- Added `Laplacian` operator. - Dockerfile support for running rocCV in a container environment. ### Changed diff --git a/docs/reference/rocCV-supported-operators.rst b/docs/reference/rocCV-supported-operators.rst index 1e214bc9..73a7fdc0 100644 --- a/docs/reference/rocCV-supported-operators.rst +++ b/docs/reference/rocCV-supported-operators.rst @@ -22,6 +22,7 @@ The rocCV is a collection of the following computer vision operators that are su "CustomCrop","Crops a region of interest from an input tensor.","U8","NHWC, HWC" "GammaContrast","Adjusts the gamma contrast on images in a tensor.","U8","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 21bb3713..6503cc7a 100644 --- a/docs/supported-operators.md +++ b/docs/supported-operators.md @@ -17,6 +17,7 @@ See below for a list of Computer Vision operators rocCV supports. |Flip|Flips the images in a tensor about the horizontal, vertical or both axes.|U8, S32, F32|NHWC, HWC|Both| |GammaContrast|Adjusts the gamma contrast of an image.|U8, U16, U32, 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| From 8f83fcec78237393e331f63e04a4d42907501c7f Mon Sep 17 00:00:00 2001 From: daniellegillai Date: Mon, 20 Jul 2026 11:51:05 -0700 Subject: [PATCH 6/7] simplified kernel struct --- include/kernels/common/laplacian_kernels.hpp | 8 ++------ include/op_laplacian.hpp | 2 +- src/op_laplacian.cpp | 1 - tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp | 6 +++--- 4 files changed, 6 insertions(+), 11 deletions(-) diff --git a/include/kernels/common/laplacian_kernels.hpp b/include/kernels/common/laplacian_kernels.hpp index 2953bcc6..4fb60766 100644 --- a/include/kernels/common/laplacian_kernels.hpp +++ b/include/kernels/common/laplacian_kernels.hpp @@ -28,18 +28,14 @@ constexpr int LaplaceKWidth = 3; constexpr int LaplaceKHeight = 3; struct LaplacianKernel { - constexpr LaplacianKernel& operator*=(float scale) { + LaplacianKernel& operator*=(float scale) { for (int i = 0; i < 9; i++) { m_kernel[i] *= scale; } return *this; } - constexpr float& operator[](int i) { - return m_kernel[i]; - } - - constexpr const float& operator[](int i) const { + __device__ __host__ const float& operator[](int i) const { return m_kernel[i]; } diff --git a/include/op_laplacian.hpp b/include/op_laplacian.hpp index 48df378e..82f9b7cd 100644 --- a/include/op_laplacian.hpp +++ b/include/op_laplacian.hpp @@ -67,7 +67,7 @@ class Laplacian final : public IOperator { * @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 to compute the second derivative filters. Must be 1 or 3. + * @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) diff --git a/src/op_laplacian.cpp b/src/op_laplacian.cpp index 9f58ce15..3d543aff 100644 --- a/src/op_laplacian.cpp +++ b/src/op_laplacian.cpp @@ -58,7 +58,6 @@ void Laplacian::operator()(hipStream_t stream, const roccv::Tensor &input, const } else if (ksize == 3) { kernel = LK3; } - if (scale != 1) { kernel *= scale; } diff --git a/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp index 9a40139f..ec1e8875 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp @@ -49,7 +49,7 @@ namespace { * @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. Must be 1 or 3. + * @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. */ @@ -99,7 +99,7 @@ std::vector GenerateGoldenLaplacian(std::vector& input, int32_t batchSiz * @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. Must be 1 or 3. + * @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. */ @@ -138,7 +138,7 @@ void TestCorrectness(int batchSize, int width, int height, ImageFormat format, i // Compare data in actual output versus the generated golden reference image // TODO check on delta - CompareVectorsNear(outputData, ref, 1); + CompareVectorsNear(outputData, ref, 1.0); } void TestNegativeLaplacian() { From 248397274d535204b7f097b0aeda72f81ef535fc Mon Sep 17 00:00:00 2001 From: daniellegillai Date: Mon, 20 Jul 2026 15:48:31 -0700 Subject: [PATCH 7/7] pybindings and refactor laplacian kernel --- CHANGELOG.md | 1 - include/op_laplacian.hpp | 2 +- .../include/operators/py_op_laplacian.hpp | 51 ++++------ python/src/main.cpp | 2 + python/src/operators/py_op_laplacian.cpp | 93 +++++++++++++++++++ src/op_laplacian.cpp | 50 ++++++++-- .../src/tests/operators/test_op_laplacian.cpp | 1 - tests/roccv/python/test_op_laplacian.py | 49 ++++++++++ 8 files changed, 201 insertions(+), 48 deletions(-) rename include/kernels/common/laplacian_kernels.hpp => python/include/operators/py_op_laplacian.hpp (61%) create mode 100644 python/src/operators/py_op_laplacian.cpp create mode 100644 tests/roccv/python/test_op_laplacian.py diff --git a/CHANGELOG.md b/CHANGELOG.md index 8c2810f7..5e89b01c 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,7 +9,6 @@ The full documentation for rocCV is available at [https://rocm.docs.amd.com/proj - Added `BrightnessContrast` operator. - Added `Gaussian` operator. - Added `Laplacian` operator. -- Added `Gaussian` operator. - Dockerfile support for running rocCV in a container environment. ### Changed diff --git a/include/op_laplacian.hpp b/include/op_laplacian.hpp index 82f9b7cd..bc46f6d0 100644 --- a/include/op_laplacian.hpp +++ b/include/op_laplacian.hpp @@ -29,7 +29,7 @@ THE SOFTWARE. namespace roccv { /** - * @brief Class for managing the Brightness Contrast operator. + * @brief Class for managing the Laplacian operator. * */ class Laplacian final : public IOperator { diff --git a/include/kernels/common/laplacian_kernels.hpp b/python/include/operators/py_op_laplacian.hpp similarity index 61% rename from include/kernels/common/laplacian_kernels.hpp rename to python/include/operators/py_op_laplacian.hpp index 4fb60766..63f9ac8c 100644 --- a/include/kernels/common/laplacian_kernels.hpp +++ b/python/include/operators/py_op_laplacian.hpp @@ -22,38 +22,19 @@ THE SOFTWARE. #pragma once -namespace Kernels { - -constexpr int LaplaceKWidth = 3; -constexpr int LaplaceKHeight = 3; - -struct LaplacianKernel { - 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]; - } - - float m_kernel[9] = {}; -}; - -// clang-format off -constexpr LaplacianKernel LK1 { - {0.0f, 1.0f, 0.0f, - 1.0f, -4.0f, 1.0f, - 0.0f, 1.0f, 0.0f} -}; - -constexpr LaplacianKernel LK3 { - {2.0f, 0.0f, 2.0f, - 0.0f, -8.0f, 0.0f, - 2.0f, 0.0f, 2.0f} -}; -// clang-format on - -} // namespace Kernels +#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 index 3d543aff..1c61639f 100644 --- a/src/op_laplacian.cpp +++ b/src/op_laplacian.cpp @@ -28,10 +28,46 @@ THE SOFTWARE. #include "common/validation_helpers.hpp" #include "core/detail/casting.hpp" #include "details/filter_2d.hpp" -#include "kernels/common/laplacian_kernels.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, +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); @@ -50,14 +86,8 @@ void Laplacian::operator()(hipStream_t stream, const roccv::Tensor &input, const eStatusType::INVALID_VALUE); } - using namespace Kernels; - LaplacianKernel kernel; - - if (ksize == 1) { - kernel = LK1; - } else if (ksize == 3) { - kernel = LK3; - } + // get the right kernel + LaplacianKernel kernel = (ksize == 1) ? LK1 : LK3; if (scale != 1) { kernel *= scale; } diff --git a/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp index ec1e8875..cb86feb5 100644 --- a/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp +++ b/tests/roccv/cpp/src/tests/operators/test_op_laplacian.cpp @@ -137,7 +137,6 @@ void TestCorrectness(int batchSize, int width, int height, ImageFormat format, i std::vector ref = GenerateGoldenLaplacian(inputData, batchSize, width, height, ksize, scale); // Compare data in actual output versus the generated golden reference image - // TODO check on delta CompareVectorsNear(outputData, ref, 1.0); } 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)