Sobel filter Kernel implementation by RooseweltMcW · Pull Request #552 · r-abishek/rpp

RooseweltMcW · 2025-12-15T17:08:19Z

No description provided.

…of case numbers (r-abishek#552) * Remove test case numbers from python scripts * Add erode and dilate test cases --------- Co-authored-by: HazarathKumarM <hazarathkumar@multicorewareinc.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>

HazarathKumarM · 2026-01-13T07:23:52Z

docs/data/doxygenOutputs/filter_augmentations_sobel_filter_kSize3_img150x150.png

only one image is added, but we have 3 types of sobel filter and different kernel sizes. check and add the missing images here

Added for all the kernel size and gradients, Done

fiona-gladwin · 2026-01-19T07:29:10Z

src/modules/tensor/cpu/kernel/sobel_filter.cpp

+        T *srcPtrChannel, *dstPtrChannel;
+        srcPtrChannel = srcPtrImage + (roi.xywhROI.xy.y * srcDescPtr->strides.hStride) + roi.xywhROI.xy.x;
+        dstPtrChannel = dstPtrImage;
+        if ((srcDescPtr->layout == RpptLayout::NCHW) && (dstDescPtr->layout == RpptLayout::NCHW) && (srcDescPtr->c == 1))


Blocking: This if is the only branch that stores to dstPtr. For every other layout/channel combination we fall through and still return RPP_SUCCESS, so NHWC tensors (which the API docs advertise) or NCHW inputs with c==3 silently produce garbage. Please add real handling for those cases, or reject them before returning success.

fiona-gladwin · 2026-01-19T07:29:25Z

src/modules/tensor/rppt_tensor_filter_augmentations.cpp

+    {
+        RpptSubpixelLayout srcSubpixelLayout = RpptSubpixelLayout::RGBtype;
+        tempPtr = rpp::deref(rppHandle).GetInitHandle()->mem.mcpu.scratchBufferHost;
+        rppt_color_to_greyscale_host(srcPtr, srcDescPtr, tempPtr, dstDescPtr, srcSubpixelLayout, rppHandle);


Blocking: rppt_color_to_greyscale_host returns an RppStatus. If that conversion fails (for example because the destination descriptor is NHWC), we plow ahead on uninitialized scratch data and later return RPP_SUCCESS. Please capture the status and propagate failures instead of ignoring them.

fiona-gladwin · 2026-01-19T07:29:35Z

src/modules/tensor/rppt_tensor_filter_augmentations.cpp

+
+    if ((srcDescPtr->dataType == RpptDataType::U8) && (dstDescPtr->dataType == RpptDataType::U8))
+    {
+        sobel_filter_host_tensor(static_cast<Rpp8u*>(tempPtr) + srcDescPtr->offsetInBytes,


Blocking: When the source tensor is already single-channel, tempPtr still aliases the original source data, but we describe it with dstDescPtr here. Any src/dst stride/layout mismatch (e.g. NHWC source going to an NCHW output) means we read the wrong pixels yet return success. Conversely, in the RGB path you now have scratch memory in the destination layout, but you are still adding srcDescPtr->offsetInBytes. Please use the descriptor/offset that actually matches tempPtr in each case.

fiona-gladwin · 2026-01-19T07:29:44Z

src/modules/tensor/rppt_tensor_filter_augmentations.cpp

+        CHECK_RETURN_STATUS(hipMalloc(&tempPtr, dataSize));
+
+        RpptSubpixelLayout srcSubpixelLayout = RpptSubpixelLayout::RGBtype;
+        rppt_color_to_greyscale_gpu(srcPtr, srcDescPtr, tempPtr, dstDescPtr, srcSubpixelLayout, rppHandle);        


Blocking: Same issue on the HIP path—the status from rppt_color_to_greyscale_gpu is ignored. If the conversion fails we still launch the Sobel kernels on whatever happens to be in tempPtr and return success. Please bail out (and free the scratch allocation) when the conversion reports an error.

…of case numbers (r-abishek#552) * Remove test case numbers from python scripts * Add erode and dilate test cases --------- Co-authored-by: HazarathKumarM <hazarathkumar@multicorewareinc.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>

fiona-gladwin · 2026-01-21T09:23:15Z

api/rppdefs.h

    /*! \brief The specified axis is invalid or out of range. (Needs to adhere to function specification.) \ingroup group_rppdefs */
-    RPP_ERROR_INVALID_AXIS              = -26
+    RPP_ERROR_INVALID_AXIS              = -26,
+    /*! \brief Invalid src tensor layout or channels. (Needs to adhere to function specification.) \ingroup group_rppdefs */


We already have INVALID_SRC_LAYOUT or INVALID_CHANNELS
Why do we need a new one?

For Sobel filter, we need both channels and layout check, so we need combined them into a single check and added a new ENUM here.
if it's not needed, we can have two separate checks

fiona-gladwin · 2026-01-21T09:23:39Z

src/include/common/cpu/rpp_cpu_batchpd.hpp

 #include <cstring>
 #include <omp.h>
-#include <half/half.hpp>
+#if __has_include(<half/half.hpp>)


Why this change is added?

Resolved, Done.

fiona-gladwin · 2026-01-21T09:24:16Z

src/include/common/op_kernel_args.hpp

 #include <type_traits>
 #include <cstdint>
-#include <half/half.hpp>
+#if __has_include(<half/half.hpp>)


Why it is added in this PR?

fiona-gladwin · 2026-01-21T09:27:53Z

src/modules/tensor/cpu/kernel/sobel_filter.cpp

+/*
+MIT License
+
+Copyright (c) 2019 - 2024 Advanced Micro Devices, Inc.


Check copyright, if this file was introduced in 2024 only modify the copyright accordingly
Same is applicable for all new files in this PR

Modified, Done.

fiona-gladwin · 2026-01-21T09:30:02Z

src/modules/tensor/cpu/kernel/sobel_filter.cpp

+#include "host_tensor_executors.hpp"
+#include "rpp_cpu_filter.hpp"
+
+Rpp32f sobel3x3X[9] = {-1, 0, 1,


Why do we have all thes variables as global?
Can we not have them as local variables, you can define it static const since it won't be modified

fiona-gladwin · 2026-01-21T09:51:40Z

utilities/test_suite/HIP/runImageTests.py

                    result = subprocess.Popen([buildFolderPath + "/build/Tensor_image_hip", srcPath1, srcPath2, dstPathTemp, str(bitDepth.value), str(outputFormatToggle.value), str(case), str(swapOrder), str(numRuns), str(testType), str(layout), "0", str(qaMode), str(decoderType), str(batchSize)] + roiList + [scriptPath], stdout=subprocess.PIPE, stderr=subprocess.PIPE)    # nosec
                    log_detected(result, errorLog, imageAugmentationMap[int(case)][0], get_bit_depth(bitDepth.value), get_image_layout_type(layout, outputFormatToggle.value, "HIP"))
+            elif imageAugmentationMap[int(case)][0] == "sobel_filter":
+                for kernelSizeAndGradient in range(9):


Why are we using range(9) here?

We have support for only 3, 5 and 7 right?

kernelSize ranges from 3 to 7 but there are 3 gradient for each kernelSize (X, Y, XY) hence 3x3 - 9 range is set as additionalParam

…ferent conditions

r-abishek · 2026-02-18T00:48:33Z

@HazarathKumarM ToT merge and Unified API changes to be reflected.

Sobel filter implementation for HOST with QA for U8 and F32

4ead48b

HazarathKumarM mentioned this pull request Dec 22, 2025

RPP Sobel Filter Tensor HOST and HIP #345

Closed

Fixed HIP F32 image generation and QA passed for 1e-4

c38f8ad

HazarathKumarM changed the base branch from develop to master January 13, 2026 07:29

HazarathKumarM changed the base branch from master to develop January 13, 2026 07:29

HazarathKumarM reviewed Jan 13, 2026

View reviewed changes

RooseweltMcW force-pushed the apr/sobel branch from 3dfcd80 to c38f8ad Compare January 13, 2026 08:54

HazarathKumarM added 4 commits January 13, 2026 04:03

Merge remote-tracking branch 'tot/develop' into apr/sobel

4ececc2

Revert test suite image file

b4401cc

Updated HOST backend implementation for sobel filter

7f3c15e

Updated documentation and added images

141fb83

fiona-gladwin reviewed Jan 19, 2026

View reviewed changes

HazarathKumarM added 2 commits January 19, 2026 05:11

Merge remote-tracking branch 'tot/develop' into apr/sobel

4751632

Resolved review comments

72032c0

fiona-gladwin reviewed Jan 21, 2026

View reviewed changes

Merge remote-tracking branch 'origin/develop' into apr/sobel

7f045a4

fiona-gladwin reviewed Jan 21, 2026

View reviewed changes

HazarathKumarM changed the base branch from develop to master January 21, 2026 09:58

HazarathKumarM changed the base branch from master to develop January 21, 2026 09:58

Resolved review comments and modified error status to set for two dif…

68d6b88

…ferent conditions

HazarathKumarM added 3 commits February 17, 2026 20:12

Merge branch 'develop' into apr/sobel

b3f08ef

updated unified api for sobel filter

3e37e72

cleanup sobel filter api

7df5225

r-abishek approved these changes Feb 18, 2026

View reviewed changes

r-abishek changed the base branch from develop to ar/opt_sobel_filter February 18, 2026 23:28

r-abishek assigned HazarathKumarM Feb 18, 2026

r-abishek added the enhancement New feature or request label Feb 18, 2026

r-abishek merged commit 07071d4 into r-abishek:ar/opt_sobel_filter Feb 18, 2026

Conversation

RooseweltMcW commented Dec 15, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

r-abishek commented Feb 18, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants