Repository navigation
Conversation
dbrockho
requested review from
AlexBrownAMD,
Jinp800125,
KKyang,
Serge45,
TonyYHsieh,
aazz44ss,
babakpst,
bnemanich,
davidd-amd,
hcman2,
imcarsonliao,
jichangjichang,
msujon-AMD,
solaslin and
vin-huang
as code owners
July 11, 2025 11:49
dbrockho
changed the title
dbrockho
changed the title
ammallya
pushed a commit
that referenced
this pull request
Oct 28, 2025
…lude Clean up build [ROCm/rocwmma commit: 74d8760]
ammallya
pushed a commit
that referenced
this pull request
Jun 8, 2026
* Updates for box filter to test * Box Filter Newer Commit with fixes * Further updates to match PKD3 and PLN3 output * Fixes for RPP Box filter border replicate - Fix accuracy * Rename variables * Restore AVX and Update AVX code for box filter * Box Filter float updates * Updates to float version * Compilation fixes * HIP Updates for Box Filter kernelSize = 3 * HIP Updates for PLN Code * HIP Updates for kernelSize = 5 * HIP Updates for kernelSize = 7 * HIP Updates for kernelSize = 9 * Remove additional else * Fix issues with alignedLength for kernelSize 3 float variants PKD3 to PLN3 * Add golden outputs * Add additional borderType parameter * Separate float implementation for 3x3 box filter planar * Rename function * Introduce functions to calculate in float type itself * Updates for PKD variants * Compilation fixes * Updates for PKD3 to PLN3 * Fix accuracy issues * PLN3 to PKD3 updates for box filter * Float shared variables * Overload box filter for various kernelSizes * Template the PKD3 and PLN3 implementations * Template the rest of the implementations * Initial cleanup of HIP Code * Cleanup of code with parameter passing launch of kernel * Add further comments * Add more comments and minor cleanups further * Initial cleanup of cpu kernel code * Add comments for padding * Add reference outputs in docs/data/doxygenOutputs * Update documentation * Minor corrections * Updates to move filter structure as common code * Update the separation based on type * Update rpp_hip_load_store.hpp * Updates for separating i8 for PLN3 code * Fixes for PLN3 and further updates for kernelSize = 3 * I8 updates for kernelSize = 5 * Updates for kernelSize = 7 * Updates for kernelSize = 9 * Fixes for kernelSize = 9 * Remove saturate pixel * Rename function * Update comment * Add declarations for I8 functions with rounding * Whitespace and type updates * Updates for performance - U8/I8 * Templated the box filter compute functions for all kernel sizes and optimized the I8 variants * Update comments * Make initial changes to template unpack function * Fixes for box filter compilation * modified padding load logic * Update the round function used * Rename verticalDirection and horizontalDirection to padVertical and padHorizontal * Fix accuracy issues --------- Co-authored-by: Abishek <52214183+r-abishek@users.noreply.github.com> Co-authored-by: HazarathKumarM <hazarathkumar@multicorewareinc.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com> Co-authored-by: Lakshmi Kumar <lakshmi.kumar@amd.com>
kliegeois
pushed a commit
to kliegeois/rocm-libraries
that referenced
this pull request
Jun 8, 2026
* Updates for box filter to test * Box Filter Newer Commit with fixes * Further updates to match PKD3 and PLN3 output * Fixes for RPP Box filter border replicate - Fix accuracy * Rename variables * Restore AVX and Update AVX code for box filter * Box Filter float updates * Updates to float version * Compilation fixes * HIP Updates for Box Filter kernelSize = 3 * HIP Updates for PLN Code * HIP Updates for kernelSize = 5 * HIP Updates for kernelSize = 7 * HIP Updates for kernelSize = 9 * Remove additional else * Fix issues with alignedLength for kernelSize 3 float variants PKD3 to PLN3 * Add golden outputs * Add additional borderType parameter * Separate float implementation for 3x3 box filter planar * Rename function * Introduce functions to calculate in float type itself * Updates for PKD variants * Compilation fixes * Updates for PKD3 to PLN3 * Fix accuracy issues * PLN3 to PKD3 updates for box filter * Float shared variables * Overload box filter for various kernelSizes * Template the PKD3 and PLN3 implementations * Template the rest of the implementations * Initial cleanup of HIP Code * Cleanup of code with parameter passing launch of kernel * Add further comments * Add more comments and minor cleanups further * Initial cleanup of cpu kernel code * Add comments for padding * Add reference outputs in docs/data/doxygenOutputs * Update documentation * Minor corrections * Updates to move filter structure as common code * Update the separation based on type * Update rpp_hip_load_store.hpp * Updates for separating i8 for PLN3 code * Fixes for PLN3 and further updates for kernelSize = 3 * I8 updates for kernelSize = 5 * Updates for kernelSize = 7 * Updates for kernelSize = 9 * Fixes for kernelSize = 9 * Remove saturate pixel * Rename function * Update comment * Add declarations for I8 functions with rounding * Whitespace and type updates * Updates for performance - U8/I8 * Templated the box filter compute functions for all kernel sizes and optimized the I8 variants * Update comments * Make initial changes to template unpack function * Fixes for box filter compilation * modified padding load logic * Update the round function used * Rename verticalDirection and horizontalDirection to padVertical and padHorizontal * Fix accuracy issues --------- Co-authored-by: Abishek <52214183+r-abishek@users.noreply.github.com> Co-authored-by: HazarathKumarM <hazarathkumar@multicorewareinc.com> Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com> Co-authored-by: Lakshmi Kumar <lakshmi.kumar@amd.com> [ROCm/rpp commit: 89de11e]
mpanoop
pushed a commit
to mpanoop/rocm-libraries
that referenced
this pull request
Oct 5, 2026
* Remove website URL from comments (ROCm#600) Referencing or using code from some websites is prohibited in this repository. This change removes an informational reference in the comments. * Fix rare memory access faults when using internal serial merge (ROCm#597) * test: add tests for internal serial merge function * refactor(detail/merge_path.hpp): removed code duplication * fix(detail/merge_path.hpp): stricter boundary checking in serial merge * fix(detail/block_sort_merge.hpp): fix missing block-wide sync During a previous refactor, serial_merge does no longer do a block sync. This has now been re-added. * feat: add unsafe variant of serial merge * fix: use bounded version for serial merge to fix rare page faults * test(test_internal_merge_path): clean up internal merge path tests * style: standardize range_t<> construction * fix(detail/merge_path.hpp): fix 'range_t<>::count1()' and 'range_t<>::count2()' return types to be same as encapsulated type * perf(detail/merge_path.hpp): use const ref in function parameters * refactor(detail/merge_path.hpp): replace redundant use of 'OffsetT' with 'unsigned int' * chore: update changelog * fix: restore missing thread sync This got removed during a rebase. * Add gfx1151 target (ROCm#601) (ROCm#603) Co-authored-by: Stanley Tsang <stanley.tsang@amd.com> * Merge back 6.2 hotfixes (ROCm#607) (ROCm#620) * Update dependency names for static builds (ROCm#557) This also removes the line setting `BUILD_SHARED_LIBS` to `ON`, which was previously required to get the correctly named packages when not specifically compiling for a static build. Updates to the ROCmCMakeBuildTools (rocm-cmake) should mean this is no longer necessary. * Fix BUILD_SHARED_LIBS for packaging (ROCm#558) * Fix the dependencies of the static packages (ROCm#563) * cmake: don't set CMAKE_C_COMPILER, as rocPRIM is a CXX project (ROCm#567) * add developer guidelines (ROCm#555) (ROCm#574) * Update Read the Docs config to Python 3.10 and latest rocm-docs-core (ROCm#564) (ROCm#579) * Cherry-pick: Optimize block_reduce_warp_reduce when block size is the same as warp size (ROCm#599) * Optimize block_reduce_warp_reduce when block size == warp size * Make conditional constexpr * Fix conflict in concepts.rst --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> * Changed precondition for edge case in serial_merge to prevent assertion error (ROCm#622) * added std::min to ensure no out of bound acess * fixed typo keys->keys1 * updated changelog * reverted std::min * implemented suggested logic * edited to conform to standards (ROCm#618) * Memory leak fix for multiple rocPRIM unit tests (ROCm#614) * fixed mem leak in test_config_dispatch.cpp * added missing hip free for method==4 in test_block_scan.kernels * added graphHelpeer class that does not cause memory leak due to using hipGraphCreate * replaced old hipGraph helpers with new class in device_bin_search * changed HIP_CHECK_NON_VOID to HIP_CHECK * fixed mem leak in device_bin_search * added additional functions * changed out old calls to hipGraphCrete to new GraphHelper class * added missing stream sync for hipgrag_algs * n * added missing hipFree and HIP_CHECK for lookback_reproducibility * added missing hipFree in test_discard_iterator * fixed test failures * removed extra hipFree * removed unused variables * updated change log * removed redundant function --------- Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> * changed to using gHelper * merged 6.3 version of merge_sort_merge_path (failing large indices test) * updated mergepath to have fix for large indices * updated the changelog for 6.3 (ROCm#632) * updated changelog for test_utils_hipgraph change * updated changelog --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nara <nara@streamhpc.com> Co-authored-by: amd-garydeng <garydeng@amd.com> Co-authored-by: Stanley Tsang <stanley.tsang@amd.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com> Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> [ROCm/rocPRIM commit: 44159bd]
mpanoop
pushed a commit
to mpanoop/rocm-libraries
that referenced
this pull request
Oct 5, 2026
* Remove website URL from comments (ROCm#600) Referencing or using code from some websites is prohibited in this repository. This change removes an informational reference in the comments. * Fix rare memory access faults when using internal serial merge (ROCm#597) * test: add tests for internal serial merge function * refactor(detail/merge_path.hpp): removed code duplication * fix(detail/merge_path.hpp): stricter boundary checking in serial merge * fix(detail/block_sort_merge.hpp): fix missing block-wide sync During a previous refactor, serial_merge does no longer do a block sync. This has now been re-added. * feat: add unsafe variant of serial merge * fix: use bounded version for serial merge to fix rare page faults * test(test_internal_merge_path): clean up internal merge path tests * style: standardize range_t<> construction * fix(detail/merge_path.hpp): fix 'range_t<>::count1()' and 'range_t<>::count2()' return types to be same as encapsulated type * perf(detail/merge_path.hpp): use const ref in function parameters * refactor(detail/merge_path.hpp): replace redundant use of 'OffsetT' with 'unsigned int' * chore: update changelog * fix: restore missing thread sync This got removed during a rebase. * Add gfx1151 target (ROCm#601) (ROCm#603) Co-authored-by: Stanley Tsang <stanley.tsang@amd.com> * Merge back 6.2 hotfixes (ROCm#607) (ROCm#620) * Update dependency names for static builds (ROCm#557) This also removes the line setting `BUILD_SHARED_LIBS` to `ON`, which was previously required to get the correctly named packages when not specifically compiling for a static build. Updates to the ROCmCMakeBuildTools (rocm-cmake) should mean this is no longer necessary. * Fix BUILD_SHARED_LIBS for packaging (ROCm#558) * Fix the dependencies of the static packages (ROCm#563) * cmake: don't set CMAKE_C_COMPILER, as rocPRIM is a CXX project (ROCm#567) * add developer guidelines (ROCm#555) (ROCm#574) * Update Read the Docs config to Python 3.10 and latest rocm-docs-core (ROCm#564) (ROCm#579) * Cherry-pick: Optimize block_reduce_warp_reduce when block size is the same as warp size (ROCm#599) * Optimize block_reduce_warp_reduce when block size == warp size * Make conditional constexpr * Fix conflict in concepts.rst --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> * Changed precondition for edge case in serial_merge to prevent assertion error (ROCm#622) * added std::min to ensure no out of bound acess * fixed typo keys->keys1 * updated changelog * reverted std::min * implemented suggested logic * edited to conform to standards (ROCm#618) * Memory leak fix for multiple rocPRIM unit tests (ROCm#614) * fixed mem leak in test_config_dispatch.cpp * added missing hip free for method==4 in test_block_scan.kernels * added graphHelpeer class that does not cause memory leak due to using hipGraphCreate * replaced old hipGraph helpers with new class in device_bin_search * changed HIP_CHECK_NON_VOID to HIP_CHECK * fixed mem leak in device_bin_search * added additional functions * changed out old calls to hipGraphCrete to new GraphHelper class * added missing stream sync for hipgrag_algs * n * added missing hipFree and HIP_CHECK for lookback_reproducibility * added missing hipFree in test_discard_iterator * fixed test failures * removed extra hipFree * removed unused variables * updated change log * removed redundant function --------- Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> * updated the changelog for 6.3 (ROCm#632) * updated default gpu to include gfx12 and gfx1151 * updated changelog * fixed minor grammar mistake in changelog * Update CHANGELOG.md Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com> * Remove gfx940,gfx941 targets (ROCm#639) * Update rocPRIM version --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nara <nara@streamhpc.com> Co-authored-by: amd-garydeng <garydeng@amd.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> Co-authored-by: Di Nguyen <dinguyennhu@gmail.com> Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com> Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> Co-authored-by: Val Movsik <160653499+vamovsik@users.noreply.github.com> [ROCm/rocPRIM commit: ffd6685]
mpanoop
pushed a commit
to mpanoop/rocm-libraries
that referenced
this pull request
Oct 5, 2026
* Remove website URL from comments (ROCm#600) Referencing or using code from some websites is prohibited in this repository. This change removes an informational reference in the comments. * Fix rare memory access faults when using internal serial merge (ROCm#597) * test: add tests for internal serial merge function * refactor(detail/merge_path.hpp): removed code duplication * fix(detail/merge_path.hpp): stricter boundary checking in serial merge * fix(detail/block_sort_merge.hpp): fix missing block-wide sync During a previous refactor, serial_merge does no longer do a block sync. This has now been re-added. * feat: add unsafe variant of serial merge * fix: use bounded version for serial merge to fix rare page faults * test(test_internal_merge_path): clean up internal merge path tests * style: standardize range_t<> construction * fix(detail/merge_path.hpp): fix 'range_t<>::count1()' and 'range_t<>::count2()' return types to be same as encapsulated type * perf(detail/merge_path.hpp): use const ref in function parameters * refactor(detail/merge_path.hpp): replace redundant use of 'OffsetT' with 'unsigned int' * chore: update changelog * fix: restore missing thread sync This got removed during a rebase. * Add gfx1151 target (ROCm#601) (ROCm#603) Co-authored-by: Stanley Tsang <stanley.tsang@amd.com> * Merge back 6.2 hotfixes (ROCm#607) (ROCm#620) * Update dependency names for static builds (ROCm#557) This also removes the line setting `BUILD_SHARED_LIBS` to `ON`, which was previously required to get the correctly named packages when not specifically compiling for a static build. Updates to the ROCmCMakeBuildTools (rocm-cmake) should mean this is no longer necessary. * Fix BUILD_SHARED_LIBS for packaging (ROCm#558) * Fix the dependencies of the static packages (ROCm#563) * cmake: don't set CMAKE_C_COMPILER, as rocPRIM is a CXX project (ROCm#567) * add developer guidelines (ROCm#555) (ROCm#574) * Update Read the Docs config to Python 3.10 and latest rocm-docs-core (ROCm#564) (ROCm#579) * Cherry-pick: Optimize block_reduce_warp_reduce when block size is the same as warp size (ROCm#599) * Optimize block_reduce_warp_reduce when block size == warp size * Make conditional constexpr * Fix conflict in concepts.rst --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> * Changed precondition for edge case in serial_merge to prevent assertion error (ROCm#622) * added std::min to ensure no out of bound acess * fixed typo keys->keys1 * updated changelog * reverted std::min * implemented suggested logic * edited to conform to standards (ROCm#618) * Memory leak fix for multiple rocPRIM unit tests (ROCm#614) * fixed mem leak in test_config_dispatch.cpp * added missing hip free for method==4 in test_block_scan.kernels * added graphHelpeer class that does not cause memory leak due to using hipGraphCreate * replaced old hipGraph helpers with new class in device_bin_search * changed HIP_CHECK_NON_VOID to HIP_CHECK * fixed mem leak in device_bin_search * added additional functions * changed out old calls to hipGraphCrete to new GraphHelper class * added missing stream sync for hipgrag_algs * n * added missing hipFree and HIP_CHECK for lookback_reproducibility * added missing hipFree in test_discard_iterator * fixed test failures * removed extra hipFree * removed unused variables * updated change log * removed redundant function --------- Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> * changed to using gHelper * merged 6.3 version of merge_sort_merge_path (failing large indices test) * updated mergepath to have fix for large indices * updated the changelog for 6.3 (ROCm#632) * updated changelog for test_utils_hipgraph change * updated changelog --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nara <nara@streamhpc.com> Co-authored-by: amd-garydeng <garydeng@amd.com> Co-authored-by: Stanley Tsang <stanley.tsang@amd.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com> Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> [ROCm/rocPRIM commit: 44159bd]
mpanoop
pushed a commit
to mpanoop/rocm-libraries
that referenced
this pull request
Oct 5, 2026
* Remove website URL from comments (ROCm#600) Referencing or using code from some websites is prohibited in this repository. This change removes an informational reference in the comments. * Fix rare memory access faults when using internal serial merge (ROCm#597) * test: add tests for internal serial merge function * refactor(detail/merge_path.hpp): removed code duplication * fix(detail/merge_path.hpp): stricter boundary checking in serial merge * fix(detail/block_sort_merge.hpp): fix missing block-wide sync During a previous refactor, serial_merge does no longer do a block sync. This has now been re-added. * feat: add unsafe variant of serial merge * fix: use bounded version for serial merge to fix rare page faults * test(test_internal_merge_path): clean up internal merge path tests * style: standardize range_t<> construction * fix(detail/merge_path.hpp): fix 'range_t<>::count1()' and 'range_t<>::count2()' return types to be same as encapsulated type * perf(detail/merge_path.hpp): use const ref in function parameters * refactor(detail/merge_path.hpp): replace redundant use of 'OffsetT' with 'unsigned int' * chore: update changelog * fix: restore missing thread sync This got removed during a rebase. * Add gfx1151 target (ROCm#601) (ROCm#603) Co-authored-by: Stanley Tsang <stanley.tsang@amd.com> * Merge back 6.2 hotfixes (ROCm#607) (ROCm#620) * Update dependency names for static builds (ROCm#557) This also removes the line setting `BUILD_SHARED_LIBS` to `ON`, which was previously required to get the correctly named packages when not specifically compiling for a static build. Updates to the ROCmCMakeBuildTools (rocm-cmake) should mean this is no longer necessary. * Fix BUILD_SHARED_LIBS for packaging (ROCm#558) * Fix the dependencies of the static packages (ROCm#563) * cmake: don't set CMAKE_C_COMPILER, as rocPRIM is a CXX project (ROCm#567) * add developer guidelines (ROCm#555) (ROCm#574) * Update Read the Docs config to Python 3.10 and latest rocm-docs-core (ROCm#564) (ROCm#579) * Cherry-pick: Optimize block_reduce_warp_reduce when block size is the same as warp size (ROCm#599) * Optimize block_reduce_warp_reduce when block size == warp size * Make conditional constexpr * Fix conflict in concepts.rst --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> * Changed precondition for edge case in serial_merge to prevent assertion error (ROCm#622) * added std::min to ensure no out of bound acess * fixed typo keys->keys1 * updated changelog * reverted std::min * implemented suggested logic * edited to conform to standards (ROCm#618) * Memory leak fix for multiple rocPRIM unit tests (ROCm#614) * fixed mem leak in test_config_dispatch.cpp * added missing hip free for method==4 in test_block_scan.kernels * added graphHelpeer class that does not cause memory leak due to using hipGraphCreate * replaced old hipGraph helpers with new class in device_bin_search * changed HIP_CHECK_NON_VOID to HIP_CHECK * fixed mem leak in device_bin_search * added additional functions * changed out old calls to hipGraphCrete to new GraphHelper class * added missing stream sync for hipgrag_algs * n * added missing hipFree and HIP_CHECK for lookback_reproducibility * added missing hipFree in test_discard_iterator * fixed test failures * removed extra hipFree * removed unused variables * updated change log * removed redundant function --------- Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> * updated the changelog for 6.3 (ROCm#632) * updated default gpu to include gfx12 and gfx1151 * updated changelog * fixed minor grammar mistake in changelog * Update CHANGELOG.md Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com> * Remove gfx940,gfx941 targets (ROCm#639) * Update rocPRIM version --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nara <nara@streamhpc.com> Co-authored-by: amd-garydeng <garydeng@amd.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Steve Leung <Steve.Leung@amd.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com> Co-authored-by: Nol Moonen <nol@streamhpc.com> Co-authored-by: Sam Wu <22262939+samjwu@users.noreply.github.com> Co-authored-by: Di Nguyen <dinguyennhu@gmail.com> Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com> Co-authored-by: Your Name <you@example.com> Co-authored-by: root <root@x1001c6s0b1n0.hostmgmt2001.cm.lockhart.amd.com> Co-authored-by: Val Movsik <160653499+vamovsik@users.noreply.github.com> [ROCm/rocPRIM commit: ffd6685]
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.
0 unused kernels removed from base logic file
1 unused kernels removed from incremental logic file
2 size(s) and 17 kernel(s) added, 0 kernel(s) removed