Repository navigation
[HipDNN][ALMIOPEN-286] - Properly enabled tidy, and fixed errors - #1820
Merged
Merged
Conversation
mousdahl-amd
requested review from
BrianHarrisonAMD,
SamuelReeder,
adickin-amd,
averinevg,
bghimireamd and
jerehartAMD
September 25, 2025 16:50
adickin-amd
changed the title
|
Github action triggered OSDB jenkins job: http://rocm-ci.amd.com/job/compute-rocm-dkms-mathlibs-osdb/379 |
assistant-librarian Bot
pushed a commit
to ROCm/hipDNN
that referenced
this pull request
Sep 25, 2025
[HipDNN][ALMIOPEN-286] - Properly enabled tidy, and fixed errors (#1820) ## Motivation clang-tidy was misconfigured and a lot of header files were getting missed. Properly enabled the clang-tidy filter, and then fixed the subsequently found errors. ## Technical Details ## Test Plan - [x] - All unit tests run - [x] - All integration tests run - [x] - ASAN run ## Test Result All tests completed successfully. There is an unrelated ASAN error in the JSON stuff. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. - [x] Code formatted
qiangpan2
pushed a commit
to qiangpan2/rocm-libraries
that referenced
this pull request
May 29, 2026
Co-authored-by: mtgu0705 <mtgu@amd.com> [ROCm/composable_kernel commit: 86d1b46]
LakshmiKumar23
added a commit
that referenced
this pull request
Oct 30, 2026
… guards (#14765) Issue Tracking JIRA ID : ROCM-33746 Fixes **Issue 1** (Blur / GaussianFilter / MedianFilter wrong on both backends) and **Issue 2** (Dilate / Erode wrong on HIP) of that ticket. ## Motivation `rppt_box_filter`, `rppt_median_filter`, `rppt_gaussian_filter`, `rppt_emboss`, `rppt_erode` and `rppt_dilate` reject any src descriptor whose `offsetInBytes` is below `12 * (kernelSize / 2)`: ```cpp if (srcDescPtr->offsetInBytes < 12 * (kernelSize / 2)) return RPP_ERROR_LOW_OFFSET; ``` The premise was that the HIP kernels load the first tile unconditionally and therefore need caller-supplied lead-in slack. That has not been true since the kernels were given their own border handling - box_filter in #597 ("Fixes for RPP Box filter border replicate"), erode/dilate in #666. Every one of these kernels now guards its fast vector load: ```cpp if ((id_x_i >= 0) && ((id_x_i + 7) < roiWidth) && (id_y_i >= 0) && (id_y_i < roiHeight)) ``` and falls back to a per-element path that clamps both coordinates into `[0, roiWidth-1] x [0, roiHeight-1]` before touching `srcPtr`. No read can leave the ROI, so no caller padding is required. The guard is vestigial. Keeping it forced callers to invent an offset they had not actually allocated, which is exactly what the OpenVX nodes in MIVisionX did: they set `offsetInBytes = 12 * (kernelSize / 2)` on an unpadded tensor, so RPP began reading 12 bytes into the image and ran 12 bytes past the end. Every filtered image came out shifted left by `offsetInBytes / channels` pixels - 4 px for RGB, 12 px for greyscale - on both the HOST and the HIP path. ## This is now urgent MIVisionX #1820 has already merged the caller-side half of this fix (`offsetInBytes = 0` in Blur, Dilate, Erode, GaussianFilter and MedianFilter). With that change in place and this guard still present, the five filter ops no longer produce a shifted image - they fail outright on the GPU: ``` ERROR: kernel org.rpp.Dilate exec failed (-1:(null)) [ERR] {output_routine} Exception thrown in the process routine: { process } ERROR: vxProcessGraph failed -1 ``` So MIVisionX develop and rocm-libraries develop are currently inconsistent, and the HIP filter/morphology ops are broken until this lands. ## Technical Details - Removed the six `RPP_ERROR_LOW_OFFSET` guards. - In `rppt_erode` / `rppt_dilate` the guard sat in an `if/else if` ahead of the backend dispatch, so the `else if` collapses back to a plain `if`. - `RPP_ERROR_LOW_OFFSET` itself is left in `rppdefs.h` - it is public API. - A non-zero `offsetInBytes` keeps working. `emboss_test.cpp` still passes a padded src on purpose to keep that path covered; only its comment is updated, since it claimed the pad was mandatory. No kernel code is changed. ## Test Plan 1. RPP QA ctest, both backends. 2. The GoogleTest correctness suite for the six affected operators. 3. The same suite again with the caller pad removed - `offsetBytes = 0` in box_filter / gaussian_filter / median_filter / erode / dilate tests - since neither existing suite covers `offsetInBytes = 0` on these kernels. This is the case that actually exercises the change. 4. rocAL's `tests/python_api/unit_tests.sh` end to end, stock libraries vs this branch, on both backends. ## Test Result **RPP QA ctest** - 100% passed, 0 failed: ``` 1/2 Test #14: rpp_qa_tests_tensor_image_host_all ... Passed 370.51 sec 2/2 Test #18: rpp_qa_tests_tensor_image_hip_all .... Passed 619.19 sec ``` **GoogleTest, padded (unchanged contract)** - `4800 tests in 6 suites: 3924 OK, 876 SKIPPED, 0 FAILED` **GoogleTest, `offsetInBytes = 0`** - identical pass counts with no caller pad at all: | Suite | padded | offset 0 | | --- | --- | --- | | BoxFilterTest | 654 OK | 654 OK | | GaussianFilterTest | 964 OK | 964 OK | | MedianFilterTest | 576 OK | 576 OK | | DilateTest | 490 OK | 490 OK | | ErodeTest | 490 OK | 490 OK | | | | **3174 OK, 0 FAILED** | **rocAL `unit_tests.sh` (RGB, this branch + MIVisionX #1820) - no regressions:** | Arm | before | after | newly passing | | --- | --- | --- | --- | | HOST | 64/74 | 66/74 | Blur, MedianFilter | | HIP | 65/74 | 68/74 | Blur, MedianFilter, Dilate | Blur, MedianFilter and Dilate are bit-exact against the shipped `GoldenOutputsTensor` references after the change; before it they matched only once shifted back by `offsetInBytes / channels` px (mean absolute error 0.0000 over the interior). rocAL's cpp_api unit tests show the same before/after on the same goldens. ## Not fixed here Two residual defects are uncovered by this change rather than caused by it - both were masked by the shift, and both reproduce with the pad still in place: - `gaussian_filter`: the outermost 1-px frame differs from the golden by up to 84. - `erode` on HIP: columns 0 and 415 differ by up to 247, while `dilate` is bit-exact. They are tracked separately and are not addressed by this PR. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
clang-tidy was misconfigured and a lot of header files were getting missed. Properly enabled the clang-tidy filter, and then fixed the subsequently found errors.
Technical Details
Test Plan
Test Result
All tests completed successfully. There is an unrelated ASAN error in the JSON stuff.
Submission Checklist