RPP Dilate on HOST and HIP by HazarathKumarM · Pull Request #554 · r-abishek/rpp

HazarathKumarM · 2025-12-22T09:26:43Z

No description provided.

remove commented code

…t function

Srihari-mcw · 2025-12-24T05:44:53Z

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

+                        blend_shuffle_max_7x7_host<7, 63, 1, 15, 127, 3>(&pxTemp[0], pxMaskPkd, blendRegisterOrder);
+                        blend_shuffle_max_7x7_host<7, 63, 1, 15, 127, 3>(&pxTemp[1], pxMaskPkd, blendRegisterOrder);
+
+                         if constexpr (std::is_same<T, Rpp8s>::value)


Remove an empty space before if

Srihari-mcw · 2025-12-24T05:59:42Z

src/modules/tensor/hip/kernel/dilate.cpp

@@ -26,287 +26,38 @@ SOFTWARE.

 // -------------------- Set 0 - dilate device helpers --------------------

-__device__ void dilate_3x3_row_hip_compute(uchar *srcPtr, d_float8 *dst_f8)
+// Templated dilate row compute function - works for any filter size (3, 5, 7, 9)


device void erode_row_hip_compute(T *srcPtr, d_float8 *dst_f8)

{

#pragma unroll for (int k = 0; k < 8; k++) { float minVal = static_cast<float>(srcPtr[k]); for (int j = 1; j < filterSize; j++) minVal = fminf(minVal, static_cast<float>(srcPtr[k + j])); dst_f8->f1[k] = fminf(dst_f8->f1[k], minVal); }

}

Modify function similar to erode

Srihari-mcw · 2025-12-24T06:01:43Z

src/modules/tensor/hip/kernel/dilate.cpp

+        dilate_row_hip_compute<7>(&src_smem[hipThreadIdx_y + 4][hipThreadIdx_x8], &sum_f8);
+        dilate_row_hip_compute<7>(&src_smem[hipThreadIdx_y + 5][hipThreadIdx_x8], &sum_f8);
+        dilate_row_hip_compute<7>(&src_smem[hipThreadIdx_y + 6][hipThreadIdx_x8], &sum_f8);
+        if constexpr (std::is_same<T, Rpp8s>::value)


Dont have any if else just have

rpp_hip_pack_float8_and_store8(dstPtr + dstIdx, &sum_f8);

Srihari-mcw · 2025-12-24T06:04:30Z

src/modules/tensor/hip/kernel/dilate.cpp

+            int clampedX = roiBeginX + max(0, min(id_x_i + i, (roiWidth - 1)));
+            int clampedIdx = (id_z * srcStridesNH.x) + (clampedY * srcStridesNH.y) + (clampedX * 3);
+
+            src_smem[hipThreadIdx_y_channel.x][hipThreadIdx_x8 + i] = srcPtr[clampedIdx];         // R


Pls indent R similar to G and B in all places across the code

Srihari-mcw · 2025-12-24T06:21:11Z

utilities/test_suite/HIP/Tensor_image_hip.cpp

                if (roiTypeSrc == RpptRoiType::LTRB)
                    convert_roi(roiTensorPtrDst, RpptRoiType::XYWH, dstDescPtr->n);
-
+                    


Remove the whitespace

Srihari-mcw · 2025-12-24T06:21:46Z

utilities/test_suite/HOST/Tensor_image_host.cpp


                // If DEBUG_MODE is set to 1 dump the outputs to csv files for debugging
-                if(DEBUG_MODE && iterCount == 0)
+                if (DEBUG_MODE && iterCount == 0)


Restore all these unnecessary changes

…nd aligned indent R.

…for single line condition

r-abishek

lgtm

Copilot

Pull request overview

This PR adds HOST backend support for the RPP (Rocm Performance Primitives) dilate morphological operation, extending the existing HIP-only implementation.

Changes:

Enabled dilate operation for both HOST and HIP backends in the test suite configuration
Added CPU implementation of dilate operation with support for multiple data types (U8, I8, F16, F32)
Implemented SIMD-optimized helper functions for efficient dilate computation across different kernel sizes (3x3, 5x5, 7x7, 9x9)

Reviewed changes

Copilot reviewed 7 out of 17 changed files in this pull request and generated 2 comments.

Show a summary per file

File	Description
utilities/test_suite/common.py	Updated dilate backend support from HIP-only to include HOST
utilities/test_suite/HOST/runImageTests.py	Added dilate to kernel size test configurations
utilities/test_suite/HOST/Tensor_image_host.cpp	Implemented dilate test case with HOST backend API call
src/modules/tensor/rppt_tensor_morphological_operations.cpp	Added rppt_dilate_host function with multi-datatype support
src/include/tensor/host_tensor_executors.hpp	Added function declarations for dilate HOST implementations
src/include/common/cpu/rpp_cpu_filter.hpp	Added SIMD helper functions and morphological operation utilities
api/rppt_tensor_morphological_operations.h	Added HOST API documentation and fixed HIP documentation typo

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Copilot · 2026-01-21T03:54:45Z

utilities/test_suite/HOST/Tensor_image_host.cpp

+
                    for (int i = 0; i < oBufferSize; i++)
                        refFile << static_cast<int>(*(outputu8 + i)) << ",";
+


The added blank lines (1870 and 1873) create inconsistent spacing around the for loop. This section appears to be unrelated to the dilate implementation and these formatting changes are unnecessary.

Suggested change

for (int i = 0; i < oBufferSize; i++)

refFile << static_cast<int>(*(outputu8 + i)) << ",";

for (int i = 0; i < oBufferSize; i++)

refFile << static_cast<int>(*(outputu8 + i)) << ",";

Copilot · 2026-01-21T03:54:45Z

src/include/common/cpu/rpp_cpu_filter.hpp

+{
+    /*  For PLN inputs                                                                          | For PKD inputs
+        pSrc[0] - [X01|X02|X03|X04|X05|X06|X07|X08], pSrc[1] - [X09|X10|X11|X12|X13|X14|X15|X16]| pSrc[0] - [R01|G01|B01|R02|G02|B02|R03|G03], pSrc[1] - [B03|R04|G04|B04|R05|G05|B05|R06],
+                  [X02|X03|X04|X05|X06|X07|X08|X09] (blend with mask [0000 0001] and permute)     pSrc[2] - [G06|B06|R07|G07|B07|R08|G08|B08], pSrc[3] - [R09|G09|B09|R10|G10|B10|R11|G11]


Missing pipe character '|' separator before 'pSrc[2]' in this comment line. All other similar comment blocks in this file consistently use '|' to separate PLN and PKD input descriptions.

sampath1117 and others added 7 commits September 5, 2024 14:54

add support for dilate in HOST backend

3f8d35d

minor fix in changelog

5d9a40f

added golden outputs

d105593

remove commented code

Merge branch 'develop' into dilate_rebased

7e12dd2

resolve build errors

08b30a3

Merge remote-tracking branch 'upstream/develop' into dilate_rebased

986b5f2

Add padding changes in HIP backend

d645a3a

HazarathKumarM mentioned this pull request Dec 22, 2025

RPP Tensor Support - Dilate on HOST #335

Closed

HazarathKumarM and others added 4 commits December 22, 2025 06:09

fix sigsev issues

42bd437

fix QA for 9x9 kernel

564a557

Add if condition for pack function and template for unpack and signex…

7b87b35

…t function

Fix the rename of preLoadRows and max Comments

a9e556b

Srihari-mcw reviewed Dec 24, 2025

View reviewed changes

Fix Fix remane of Loader and MorphVecLoader

e58f3a6

Srihari-mcw reviewed Dec 24, 2025

View reviewed changes

mukeshj0606 added 5 commits December 24, 2025 01:24

Fix empty space, dilate_row_hip_compute function, removed if & else a…

0d4e930

…nd aligned indent R.

Fix remove whitespace and restored all unnecessary changes.

c526b60

Fix remove precision line and reverted back to static cast.

ffee28d

Fix remove empty line, rename of kernelSze & padPolicy and remove {} …

806ae85

…for single line condition

Fix Indentation of IF condition.

f7fae26

r-abishek approved these changes Jan 21, 2026

View reviewed changes

r-abishek changed the base branch from develop to ar/opt_dilate January 21, 2026 03:53

r-abishek assigned HazarathKumarM Jan 21, 2026

r-abishek added the enhancement New feature or request label Jan 21, 2026

r-abishek requested a review from Copilot January 21, 2026 03:54

Copilot AI reviewed Jan 21, 2026

View reviewed changes

resolved review comments

e2ecd4d

r-abishek merged commit ad42d02 into r-abishek:ar/opt_dilate Jan 27, 2026

		if (roiTypeSrc == RpptRoiType::LTRB)
		convert_roi(roiTensorPtrDst, RpptRoiType::XYWH, dstDescPtr->n);


		for (int i = 0; i < oBufferSize; i++)
		refFile << static_cast<int>(*(outputu8 + i)) << ",";

Conversation

HazarathKumarM commented Dec 22, 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!

r-abishek left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Copilot AI Jan 21, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Copilot AI Jan 21, 2026

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants