Repository navigation
Conversation
…Cm#8282) ## Motivation When using the `invoke` workflow (`invoke rocisa` → edit rocisa C++ → `invoke build-client` → `Tensile/bin/Tensile`), rocisa bindings could become stale: 1. **Stale bindings bug:** `invoke build-client` only rebuilds a bundled `_rocisa.so` in `build_tmp/`, but `Tensile/bin/Tensile` imports from the *editable* `_rocisa.so` in site-packages. Editing rocisa C++ and running only `build-client` left the editable bindings stale — surfacing as an `ImportError` if `_build_info.py` was present, or silently using wrong bindings otherwise. 2. **Wasted compilation:** The bundled `_rocisa.so` in `build_tmp/` was always built (due to a flag bug) even though it's unused in the `tensilelite` preset configuration (`HIPBLASLT_ENABLE_DEVICE=OFF` means no device-library codegen). ## Technical Details ### Commit 1: Fix `--bundle-python-deps` flag The `bundle_python_deps` parameter defaulted to `False`, but the code only passed `-DHIPBLASLT_BUNDLE_PYTHON_DEPS=ON` when `True` and nothing when `False`. Since the CMake option defaults to `ON`, the flag was effectively a no-op. Fix: Always pass the value explicitly so `False` → `=OFF`. This eliminates the unnecessary compilation of the bundled `_rocisa.so` during `invoke build-client`. ### Commit 2: Refresh editable rocisa in `build-client` `build-client` now re-runs the editable install (`pip install -e rocisa/`) before building the client. This is safe and fast because: - **Conditional:** Only triggers when rocisa is installed *editable* (detected via PEP 610 `direct_url.json`). Never clobbers tox's non-editable install or attempts to install rocisa where absent. - **Graceful degradation:** Warns (doesn't fail) when build backend (`scikit-build-core`/`nanobind`) is unavailable. - **Incremental:** New persistent `build-dir` in `pyproject.toml` makes reinstalls a cmake/make no-op when nothing changed. New flag `--no-rebuild-rocisa` to opt out. ### Commit 3: Add rocisa build dependencies to tox.ini Add `scikit-build-core` to tox deps (alongside existing `nanobind`) so the rocisa build toolchain is available in tox environments. ## Test Plan - [x] `invoke rocisa` still works (initial editable install) - [x] `invoke build-client` works when rocisa is NOT installed (no-op rebuild, no failure) - [x] `invoke build-client` works when rocisa IS installed editable (triggers rebuild) - [x] After editing a rocisa `.cpp` file, `invoke build-client` followed by `import rocisa` works (no stale `ImportError`) - [x] `invoke build-client --no-rebuild-rocisa` skips the rocisa rebuild - [x] `invoke build-client --bundle-python-deps` still builds the bundled `_rocisa.so` when explicitly requested - [x] Default `invoke build-client` passes `HIPBLASLT_BUNDLE_PYTHON_DEPS=OFF` (verified in CMakeCache.txt) - [x] tox workflows (`tox -e rocisa`: 48 tests, 0 failures) ## Test Results | Test | Result | |------|--------| | `invoke rocisa` (initial editable install) | ✅ Pass | | `invoke build-client` when rocisa NOT installed | ✅ Pass (no-op, no failure) | | `invoke build-client` when rocisa IS installed editable | ✅ Pass (triggers rebuild) | | Edit `.cpp` → `invoke build-client` → `import rocisa` | ✅ Pass (no stale ImportError) | | `--no-rebuild-rocisa` skips rebuild | ✅ Pass | | `--bundle-python-deps` enables bundling | ✅ Pass (`BUNDLE=ON` in cache) | | Default passes `BUNDLE=OFF` | ✅ Pass (no wasted `_rocisa` compilation) | | Full cmake configure | ✅ Pass | | `tox -e rocisa` | ✅ Pass (48 tests, 0 failures, 0 errors) | ## Checklist - [x] Code compiles without errors - [x] New helper functions have docstrings explaining behavior and edge cases - [x] Changes are backwards compatible (new behavior is opt-out via `--no-rebuild-rocisa`) - [x] Tested on a system with rocisa build toolchain - [x] tox workflows verified Co-authored-by: Nathan Henderson <nathan.henderson@amd.com> --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…y for MGX" (ROCm#8333) Reverts ROCm#6086 Need to revert as the codegen test for fmha is failing due to including std header: 2026-06-11T22:36:03.673Z] In file included from /tmp/comgr-953928-0-473822/include/ck/host/device_fmha_fwd/fmha_fwd_wrapper.hpp:8: [2026-06-11T22:36:03.673Z] In file included from /bin/../lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/cmath:49: [2026-06-11T22:36:03.673Z] In file included from /bin/../lib/gcc/x86_64-linux-gnu/13/../../../../include/c++/13/bits/std_abs.h:38: [2026-06-11T22:36:03.673Z] /usr/include/stdlib.h:32:10: fatal error: 'stddef.h' file not found [2026-06-11T22:36:03.673Z] 32 | #include <stddef.h> [2026-06-11T22:36:03.673Z] | ^~~~~~~~~~ The ck_tile headers were never prepped for hiprtc compilation.
…inx in the tensile-docs-dependencies group (ROCm#8364) Bumps the tensile-docs-dependencies group in /shared/tensile/docs/sphinx with 1 update: [rocm-docs-core](https://github.com/ROCm/rocm-docs-core). Updates `rocm-docs-core` from 1.33.1 to 1.35.0 <details> <summary>Release notes</summary> <p><em>Sourced from <a href="https://github.com/ROCm/rocm-docs-core/releases">rocm-docs-core's releases</a>.</em></p> <blockquote> <h2>v1.35.0 (2026-06-10)</h2> <h3>Feat</h3> <ul> <li>recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li>Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li>Add llms-full.txt generation and LLM friendly guide</li> <li>Ignore .cline_storage and .vscode directories</li> </ul> <h3>Fix</h3> <ul> <li>add new repo sync pair</li> </ul> <p>[main b623885] bump: version 1.34.0 → 1.35.0 3 files changed, 17 insertions(+), 4 deletions(-)</p> <h2>v1.34.0 (2026-04-28)</h2> <h3>Feat</h3> <ul> <li>add hipFile</li> <li>Add CUID project</li> </ul> <h3>Fix</h3> <ul> <li>fix typo</li> </ul> <p>[main 98ca6ee] bump: version 1.33.1 → 1.34.0 3 files changed, 15 insertions(+), 4 deletions(-)</p> </blockquote> </details> <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md">rocm-docs-core's changelog</a>.</em></p> <blockquote> <h2>v1.35.0 (2026-06-10)</h2> <h3>Feat</h3> <ul> <li>recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li>Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li>Add llms-full.txt generation and LLM friendly guide</li> <li>Ignore .cline_storage and .vscode directories</li> </ul> <h3>Fix</h3> <ul> <li>add new repo sync pair</li> </ul> <h2>v1.34.0 (2026-04-28)</h2> <h3>Feat</h3> <ul> <li>add hipFile</li> <li>Add CUID project</li> </ul> <h3>Fix</h3> <ul> <li>fix typo</li> </ul> </blockquote> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/b623885e7e4fe2f87bd2898d8ea0e2c4ded2eca1"><code>b623885</code></a> bump: version 1.34.0 → 1.35.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/d237fdefe1b60214be99e05e0be78e04b09684be"><code>d237fde</code></a> Merge pull request <a href="https://redirect.github.com/ROCm/rocm-docs-core/issues/1543">#1543</a> from ROCm/intersphinx-preview-version-pattern</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/a2eb21e60b3ff8fc22dfec3ea0d5261ea7af0444"><code>a2eb21e</code></a> Merge pull request <a href="https://redirect.github.com/ROCm/rocm-docs-core/issues/1531">#1531</a> from ROCm/llms_full_generation</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/1d09ac5a83f44f81abd9fba6fbac13deb1a57855"><code>1d09ac5</code></a> build: bump myst-nb from 1.3.0 to 1.4.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/92aa649c8df2450600a0ffb694b2005bc9e06c35"><code>92aa649</code></a> build: bump gitpython from 3.1.45 to 3.1.50</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/d06450c2a2ea0a2adb87f9d2c6284c0f5e7bb6ec"><code>d06450c</code></a> build: bump pyjwt from 2.10.1 to 2.13.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/bdb2eeb4e9d1c343cd6c2ede949a1d480ff87ffd"><code>bdb2eeb</code></a> feat: Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/45d3e9ac19b5d55542c9fde28246767f8cab1fd7"><code>45d3e9a</code></a> feat: Add llms-full.txt generation and LLM friendly guide</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/eabde61dfe27908776c13f08670cb3fa69929ec2"><code>eabde61</code></a> feat: recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/000057c2d597ae14e3997738e52b7ca99011b2a9"><code>000057c</code></a> build: bump idna from 3.11 to 3.17</li> <li>Additional commits viewable in <a href="https://github.com/ROCm/rocm-docs-core/compare/v1.33.1...v1.35.0">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore <dependency name> major version` will close this group update PR and stop Dependabot creating any more for the specific dependency's major version (unless you unignore this specific dependency's major version or upgrade to it yourself) - `@dependabot ignore <dependency name> minor version` will close this group update PR and stop Dependabot creating any more for the specific dependency's minor version (unless you unignore this specific dependency's minor version or upgrade to it yourself) - `@dependabot ignore <dependency name>` will close this group update PR and stop Dependabot creating any more for the specific dependency (unless you unignore this specific dependency or upgrade to it yourself) - `@dependabot unignore <dependency name>` will remove all of the ignore conditions of the specified dependency - `@dependabot unignore <dependency name> <ignore condition>` will remove the ignore condition of the specified dependency and ignore conditions </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…phinx in the rocblas-docs-dependencies group (ROCm#8360) Bumps the rocblas-docs-dependencies group in /projects/rocblas/docs/sphinx with 1 update: [rocm-docs-core](https://github.com/ROCm/rocm-docs-core). Updates `rocm-docs-core` from 1.33.1 to 1.35.0 <details> <summary>Release notes</summary> <p><em>Sourced from <a href="https://github.com/ROCm/rocm-docs-core/releases">rocm-docs-core's releases</a>.</em></p> <blockquote> <h2>v1.35.0 (2026-06-10)</h2> <h3>Feat</h3> <ul> <li>recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li>Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li>Add llms-full.txt generation and LLM friendly guide</li> <li>Ignore .cline_storage and .vscode directories</li> </ul> <h3>Fix</h3> <ul> <li>add new repo sync pair</li> </ul> <p>[main b623885] bump: version 1.34.0 → 1.35.0 3 files changed, 17 insertions(+), 4 deletions(-)</p> <h2>v1.34.0 (2026-04-28)</h2> <h3>Feat</h3> <ul> <li>add hipFile</li> <li>Add CUID project</li> </ul> <h3>Fix</h3> <ul> <li>fix typo</li> </ul> <p>[main 98ca6ee] bump: version 1.33.1 → 1.34.0 3 files changed, 15 insertions(+), 4 deletions(-)</p> </blockquote> </details> <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md">rocm-docs-core's changelog</a>.</em></p> <blockquote> <h2>v1.35.0 (2026-06-10)</h2> <h3>Feat</h3> <ul> <li>recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li>Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li>Add llms-full.txt generation and LLM friendly guide</li> <li>Ignore .cline_storage and .vscode directories</li> </ul> <h3>Fix</h3> <ul> <li>add new repo sync pair</li> </ul> <h2>v1.34.0 (2026-04-28)</h2> <h3>Feat</h3> <ul> <li>add hipFile</li> <li>Add CUID project</li> </ul> <h3>Fix</h3> <ul> <li>fix typo</li> </ul> </blockquote> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/b623885e7e4fe2f87bd2898d8ea0e2c4ded2eca1"><code>b623885</code></a> bump: version 1.34.0 → 1.35.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/d237fdefe1b60214be99e05e0be78e04b09684be"><code>d237fde</code></a> Merge pull request <a href="https://redirect.github.com/ROCm/rocm-docs-core/issues/1543">#1543</a> from ROCm/intersphinx-preview-version-pattern</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/a2eb21e60b3ff8fc22dfec3ea0d5261ea7af0444"><code>a2eb21e</code></a> Merge pull request <a href="https://redirect.github.com/ROCm/rocm-docs-core/issues/1531">#1531</a> from ROCm/llms_full_generation</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/1d09ac5a83f44f81abd9fba6fbac13deb1a57855"><code>1d09ac5</code></a> build: bump myst-nb from 1.3.0 to 1.4.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/92aa649c8df2450600a0ffb694b2005bc9e06c35"><code>92aa649</code></a> build: bump gitpython from 3.1.45 to 3.1.50</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/d06450c2a2ea0a2adb87f9d2c6284c0f5e7bb6ec"><code>d06450c</code></a> build: bump pyjwt from 2.10.1 to 2.13.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/bdb2eeb4e9d1c343cd6c2ede949a1d480ff87ffd"><code>bdb2eeb</code></a> feat: Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/45d3e9ac19b5d55542c9fde28246767f8cab1fd7"><code>45d3e9a</code></a> feat: Add llms-full.txt generation and LLM friendly guide</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/eabde61dfe27908776c13f08670cb3fa69929ec2"><code>eabde61</code></a> feat: recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/000057c2d597ae14e3997738e52b7ca99011b2a9"><code>000057c</code></a> build: bump idna from 3.11 to 3.17</li> <li>Additional commits viewable in <a href="https://github.com/ROCm/rocm-docs-core/compare/v1.33.1...v1.35.0">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore <dependency name> major version` will close this group update PR and stop Dependabot creating any more for the specific dependency's major version (unless you unignore this specific dependency's major version or upgrade to it yourself) - `@dependabot ignore <dependency name> minor version` will close this group update PR and stop Dependabot creating any more for the specific dependency's minor version (unless you unignore this specific dependency's minor version or upgrade to it yourself) - `@dependabot ignore <dependency name>` will close this group update PR and stop Dependabot creating any more for the specific dependency (unless you unignore this specific dependency or upgrade to it yourself) - `@dependabot unignore <dependency name>` will remove all of the ignore conditions of the specified dependency - `@dependabot unignore <dependency name> <ignore condition>` will remove the ignore condition of the specified dependency and ignore conditions </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…phinx in the hipblas-docs-dependencies group (ROCm#8358) Bumps the hipblas-docs-dependencies group in /projects/hipblas/docs/sphinx with 1 update: [rocm-docs-core](https://github.com/ROCm/rocm-docs-core). Updates `rocm-docs-core` from 1.33.1 to 1.35.0 <details> <summary>Release notes</summary> <p><em>Sourced from <a href="https://github.com/ROCm/rocm-docs-core/releases">rocm-docs-core's releases</a>.</em></p> <blockquote> <h2>v1.35.0 (2026-06-10)</h2> <h3>Feat</h3> <ul> <li>recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li>Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li>Add llms-full.txt generation and LLM friendly guide</li> <li>Ignore .cline_storage and .vscode directories</li> </ul> <h3>Fix</h3> <ul> <li>add new repo sync pair</li> </ul> <p>[main b623885] bump: version 1.34.0 → 1.35.0 3 files changed, 17 insertions(+), 4 deletions(-)</p> <h2>v1.34.0 (2026-04-28)</h2> <h3>Feat</h3> <ul> <li>add hipFile</li> <li>Add CUID project</li> </ul> <h3>Fix</h3> <ul> <li>fix typo</li> </ul> <p>[main 98ca6ee] bump: version 1.33.1 → 1.34.0 3 files changed, 15 insertions(+), 4 deletions(-)</p> </blockquote> </details> <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md">rocm-docs-core's changelog</a>.</em></p> <blockquote> <h2>v1.35.0 (2026-06-10)</h2> <h3>Feat</h3> <ul> <li>recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li>Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li>Add llms-full.txt generation and LLM friendly guide</li> <li>Ignore .cline_storage and .vscode directories</li> </ul> <h3>Fix</h3> <ul> <li>add new repo sync pair</li> </ul> <h2>v1.34.0 (2026-04-28)</h2> <h3>Feat</h3> <ul> <li>add hipFile</li> <li>Add CUID project</li> </ul> <h3>Fix</h3> <ul> <li>fix typo</li> </ul> </blockquote> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/b623885e7e4fe2f87bd2898d8ea0e2c4ded2eca1"><code>b623885</code></a> bump: version 1.34.0 → 1.35.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/d237fdefe1b60214be99e05e0be78e04b09684be"><code>d237fde</code></a> Merge pull request <a href="https://redirect.github.com/ROCm/rocm-docs-core/issues/1543">#1543</a> from ROCm/intersphinx-preview-version-pattern</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/a2eb21e60b3ff8fc22dfec3ea0d5261ea7af0444"><code>a2eb21e</code></a> Merge pull request <a href="https://redirect.github.com/ROCm/rocm-docs-core/issues/1531">#1531</a> from ROCm/llms_full_generation</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/1d09ac5a83f44f81abd9fba6fbac13deb1a57855"><code>1d09ac5</code></a> build: bump myst-nb from 1.3.0 to 1.4.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/92aa649c8df2450600a0ffb694b2005bc9e06c35"><code>92aa649</code></a> build: bump gitpython from 3.1.45 to 3.1.50</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/d06450c2a2ea0a2adb87f9d2c6284c0f5e7bb6ec"><code>d06450c</code></a> build: bump pyjwt from 2.10.1 to 2.13.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/bdb2eeb4e9d1c343cd6c2ede949a1d480ff87ffd"><code>bdb2eeb</code></a> feat: Fix gate/emission mismatch and improve headings in llms-full.txt</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/45d3e9ac19b5d55542c9fde28246767f8cab1fd7"><code>45d3e9a</code></a> feat: Add llms-full.txt generation and LLM friendly guide</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/eabde61dfe27908776c13f08670cb3fa69929ec2"><code>eabde61</code></a> feat: recognize X.Y.Z-preview slugs as static intersphinx versions</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/000057c2d597ae14e3997738e52b7ca99011b2a9"><code>000057c</code></a> build: bump idna from 3.11 to 3.17</li> <li>Additional commits viewable in <a href="https://github.com/ROCm/rocm-docs-core/compare/v1.33.1...v1.35.0">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore <dependency name> major version` will close this group update PR and stop Dependabot creating any more for the specific dependency's major version (unless you unignore this specific dependency's major version or upgrade to it yourself) - `@dependabot ignore <dependency name> minor version` will close this group update PR and stop Dependabot creating any more for the specific dependency's minor version (unless you unignore this specific dependency's minor version or upgrade to it yourself) - `@dependabot ignore <dependency name>` will close this group update PR and stop Dependabot creating any more for the specific dependency (unless you unignore this specific dependency or upgrade to it yourself) - `@dependabot unignore <dependency name>` will remove all of the ignore conditions of the specified dependency - `@dependabot unignore <dependency name> <ignore condition>` will remove the ignore condition of the specified dependency and ignore conditions </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…Cm#8386) Bumps [tornado](https://github.com/tornadoweb/tornado) from 6.5.5 to 6.5.6. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/tornadoweb/tornado/blob/master/docs/releases.rst">tornado's changelog</a>.</em></p> <blockquote> <h1>Release notes</h1> <p>.. toctree:: :maxdepth: 2</p> <p>releases/v6.5.7 releases/v6.5.6 releases/v6.5.5 releases/v6.5.4 releases/v6.5.3 releases/v6.5.2 releases/v6.5.1 releases/v6.5.0 releases/v6.4.2 releases/v6.4.1 releases/v6.4.0 releases/v6.3.3 releases/v6.3.2 releases/v6.3.1 releases/v6.3.0 releases/v6.2.0 releases/v6.1.0 releases/v6.0.4 releases/v6.0.3 releases/v6.0.2 releases/v6.0.1 releases/v6.0.0 releases/v5.1.1 releases/v5.1.0 releases/v5.0.2 releases/v5.0.1 releases/v5.0.0 releases/v4.5.3 releases/v4.5.2 releases/v4.5.1 releases/v4.5.0 releases/v4.4.3 releases/v4.4.2 releases/v4.4.1 releases/v4.4.0 releases/v4.3.0 releases/v4.2.1 releases/v4.2.0 releases/v4.1.0 releases/v4.0.2 releases/v4.0.1 releases/v4.0.0 releases/v3.2.2 releases/v3.2.1</p> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/tornadoweb/tornado/commit/aba2569f7ed7a6bdbef816658fb6b7182531b751"><code>aba2569</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3626">#3626</a> from bdarnell/fixes-656</li> <li><a href="https://github.com/tornadoweb/tornado/commit/a24b260e0d22fd48acea1a2635526c1700e7ac09"><code>a24b260</code></a> httpclient_test: Accept an additional error message variant</li> <li><a href="https://github.com/tornadoweb/tornado/commit/a74240a70268fe5cb40a127951cb21549ab9ff24"><code>a74240a</code></a> Release notes and version bump for 6.5.6.</li> <li><a href="https://github.com/tornadoweb/tornado/commit/e8fc7edb238f1022e39f9d0b9d297fc7c21fb0a5"><code>e8fc7ed</code></a> simple_httpclient: Strip auth headers on cross-origin redirects</li> <li><a href="https://github.com/tornadoweb/tornado/commit/96dc88c2a05705287856b2cd6b4b4034f9a6aaac"><code>96dc88c</code></a> speedups: validate mask length</li> <li><a href="https://github.com/tornadoweb/tornado/commit/ff808b33adc52d89a549376a5e3628e92abbc8ff"><code>ff808b3</code></a> http1connection: Enforce max_body_size in _GzipMessageDelegate</li> <li><a href="https://github.com/tornadoweb/tornado/commit/ede4e37f93c1edbc0bf749e9a57c9db2501cd54b"><code>ede4e37</code></a> auth: Correctly parse check_authentication response</li> <li><a href="https://github.com/tornadoweb/tornado/commit/1c178bef88bbd29907eb94a2a649a4a6675681de"><code>1c178be</code></a> Remove obsolete curl force_timeout workaround</li> <li><a href="https://github.com/tornadoweb/tornado/commit/c99d55bb6cc0c9da2c6696545ed4ee1d20b7fcf0"><code>c99d55b</code></a> Replace deprecated pycurl IOCTLFUNCTION callback with SEEKFUNCTION</li> <li><a href="https://github.com/tornadoweb/tornado/commit/27614316ef8ad125fe18725cf96e384560ba0e14"><code>2761431</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3587">#3587</a> from bdarnell/fix-link</li> <li>Additional commits viewable in <a href="https://github.com/tornadoweb/tornado/compare/v6.5.5...v6.5.6">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) You can disable automated security fix PRs for this repo from the [Security Alerts page](https://github.com/ROCm/rocm-libraries/network/alerts). </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…#8385) Bumps [tornado](https://github.com/tornadoweb/tornado) from 6.5.5 to 6.5.6. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/tornadoweb/tornado/blob/master/docs/releases.rst">tornado's changelog</a>.</em></p> <blockquote> <h1>Release notes</h1> <p>.. toctree:: :maxdepth: 2</p> <p>releases/v6.5.7 releases/v6.5.6 releases/v6.5.5 releases/v6.5.4 releases/v6.5.3 releases/v6.5.2 releases/v6.5.1 releases/v6.5.0 releases/v6.4.2 releases/v6.4.1 releases/v6.4.0 releases/v6.3.3 releases/v6.3.2 releases/v6.3.1 releases/v6.3.0 releases/v6.2.0 releases/v6.1.0 releases/v6.0.4 releases/v6.0.3 releases/v6.0.2 releases/v6.0.1 releases/v6.0.0 releases/v5.1.1 releases/v5.1.0 releases/v5.0.2 releases/v5.0.1 releases/v5.0.0 releases/v4.5.3 releases/v4.5.2 releases/v4.5.1 releases/v4.5.0 releases/v4.4.3 releases/v4.4.2 releases/v4.4.1 releases/v4.4.0 releases/v4.3.0 releases/v4.2.1 releases/v4.2.0 releases/v4.1.0 releases/v4.0.2 releases/v4.0.1 releases/v4.0.0 releases/v3.2.2 releases/v3.2.1</p> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/tornadoweb/tornado/commit/aba2569f7ed7a6bdbef816658fb6b7182531b751"><code>aba2569</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3626">#3626</a> from bdarnell/fixes-656</li> <li><a href="https://github.com/tornadoweb/tornado/commit/a24b260e0d22fd48acea1a2635526c1700e7ac09"><code>a24b260</code></a> httpclient_test: Accept an additional error message variant</li> <li><a href="https://github.com/tornadoweb/tornado/commit/a74240a70268fe5cb40a127951cb21549ab9ff24"><code>a74240a</code></a> Release notes and version bump for 6.5.6.</li> <li><a href="https://github.com/tornadoweb/tornado/commit/e8fc7edb238f1022e39f9d0b9d297fc7c21fb0a5"><code>e8fc7ed</code></a> simple_httpclient: Strip auth headers on cross-origin redirects</li> <li><a href="https://github.com/tornadoweb/tornado/commit/96dc88c2a05705287856b2cd6b4b4034f9a6aaac"><code>96dc88c</code></a> speedups: validate mask length</li> <li><a href="https://github.com/tornadoweb/tornado/commit/ff808b33adc52d89a549376a5e3628e92abbc8ff"><code>ff808b3</code></a> http1connection: Enforce max_body_size in _GzipMessageDelegate</li> <li><a href="https://github.com/tornadoweb/tornado/commit/ede4e37f93c1edbc0bf749e9a57c9db2501cd54b"><code>ede4e37</code></a> auth: Correctly parse check_authentication response</li> <li><a href="https://github.com/tornadoweb/tornado/commit/1c178bef88bbd29907eb94a2a649a4a6675681de"><code>1c178be</code></a> Remove obsolete curl force_timeout workaround</li> <li><a href="https://github.com/tornadoweb/tornado/commit/c99d55bb6cc0c9da2c6696545ed4ee1d20b7fcf0"><code>c99d55b</code></a> Replace deprecated pycurl IOCTLFUNCTION callback with SEEKFUNCTION</li> <li><a href="https://github.com/tornadoweb/tornado/commit/27614316ef8ad125fe18725cf96e384560ba0e14"><code>2761431</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3587">#3587</a> from bdarnell/fix-link</li> <li>Additional commits viewable in <a href="https://github.com/tornadoweb/tornado/compare/v6.5.5...v6.5.6">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) You can disable automated security fix PRs for this repo from the [Security Alerts page](https://github.com/ROCm/rocm-libraries/network/alerts). </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
## Motivation At the time when we introduced the smart test filter to only build and run tests affected by the PR changes, we disabled the client examples, which required full CK build, and also the hiprtc tests that were grouped with the client examples. This caused a few PRs to sneak through that caused the hiprtc compilation to fail. By restoring the hiprtc tests in all PRs, we should close this gap. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
ROCm#8006) Commit ROCm#7644 ("Exclude F6 from data initialization") renamed TensileLite::Client::isMXProblem to isMXProblemExceptF6 across the header and client sources, but DataInit_test.cpp was not updated, breaking the test build with "use of undeclared identifier 'isMXProblem'". Update the using-declaration, all call sites, and the Section 2 contract comment to the new name. isMXProblemExceptF6 (not the compiler-suggested isMXFP4Problem) preserves the original test semantics, since the suite expects FP8/BFloat8/mixed MX problems to be true as well as FP4.
…OCm#8390) Bumps [tornado](https://github.com/tornadoweb/tornado) from 6.5.5 to 6.5.6. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/tornadoweb/tornado/blob/master/docs/releases.rst">tornado's changelog</a>.</em></p> <blockquote> <h1>Release notes</h1> <p>.. toctree:: :maxdepth: 2</p> <p>releases/v6.5.7 releases/v6.5.6 releases/v6.5.5 releases/v6.5.4 releases/v6.5.3 releases/v6.5.2 releases/v6.5.1 releases/v6.5.0 releases/v6.4.2 releases/v6.4.1 releases/v6.4.0 releases/v6.3.3 releases/v6.3.2 releases/v6.3.1 releases/v6.3.0 releases/v6.2.0 releases/v6.1.0 releases/v6.0.4 releases/v6.0.3 releases/v6.0.2 releases/v6.0.1 releases/v6.0.0 releases/v5.1.1 releases/v5.1.0 releases/v5.0.2 releases/v5.0.1 releases/v5.0.0 releases/v4.5.3 releases/v4.5.2 releases/v4.5.1 releases/v4.5.0 releases/v4.4.3 releases/v4.4.2 releases/v4.4.1 releases/v4.4.0 releases/v4.3.0 releases/v4.2.1 releases/v4.2.0 releases/v4.1.0 releases/v4.0.2 releases/v4.0.1 releases/v4.0.0 releases/v3.2.2 releases/v3.2.1</p> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/tornadoweb/tornado/commit/aba2569f7ed7a6bdbef816658fb6b7182531b751"><code>aba2569</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3626">#3626</a> from bdarnell/fixes-656</li> <li><a href="https://github.com/tornadoweb/tornado/commit/a24b260e0d22fd48acea1a2635526c1700e7ac09"><code>a24b260</code></a> httpclient_test: Accept an additional error message variant</li> <li><a href="https://github.com/tornadoweb/tornado/commit/a74240a70268fe5cb40a127951cb21549ab9ff24"><code>a74240a</code></a> Release notes and version bump for 6.5.6.</li> <li><a href="https://github.com/tornadoweb/tornado/commit/e8fc7edb238f1022e39f9d0b9d297fc7c21fb0a5"><code>e8fc7ed</code></a> simple_httpclient: Strip auth headers on cross-origin redirects</li> <li><a href="https://github.com/tornadoweb/tornado/commit/96dc88c2a05705287856b2cd6b4b4034f9a6aaac"><code>96dc88c</code></a> speedups: validate mask length</li> <li><a href="https://github.com/tornadoweb/tornado/commit/ff808b33adc52d89a549376a5e3628e92abbc8ff"><code>ff808b3</code></a> http1connection: Enforce max_body_size in _GzipMessageDelegate</li> <li><a href="https://github.com/tornadoweb/tornado/commit/ede4e37f93c1edbc0bf749e9a57c9db2501cd54b"><code>ede4e37</code></a> auth: Correctly parse check_authentication response</li> <li><a href="https://github.com/tornadoweb/tornado/commit/1c178bef88bbd29907eb94a2a649a4a6675681de"><code>1c178be</code></a> Remove obsolete curl force_timeout workaround</li> <li><a href="https://github.com/tornadoweb/tornado/commit/c99d55bb6cc0c9da2c6696545ed4ee1d20b7fcf0"><code>c99d55b</code></a> Replace deprecated pycurl IOCTLFUNCTION callback with SEEKFUNCTION</li> <li><a href="https://github.com/tornadoweb/tornado/commit/27614316ef8ad125fe18725cf96e384560ba0e14"><code>2761431</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3587">#3587</a> from bdarnell/fix-link</li> <li>Additional commits viewable in <a href="https://github.com/tornadoweb/tornado/compare/v6.5.5...v6.5.6">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) You can disable automated security fix PRs for this repo from the [Security Alerts page](https://github.com/ROCm/rocm-libraries/network/alerts). </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…m#8378) ## Motivation The CK Groovy library is growing and will be reorganized into a self-describing `groovy/` folder rather than living under `src/` and `vars/`. This PR creates that folder pre-emptively and adds it to the TheRock CI skip-list so that future Groovy additions do not unnecessarily trigger TheRock builds. ## Technical Details - Added `projects/composablekernel/groovy/` with a `.gitkeep` to establish the directory in the repo. - Added `"projects/composablekernel/groovy/*"` to `SKIPPABLE_PATH_PATTERNS` in `.github/scripts/therock_configure_ci.py` alongside the existing `vars/*` entry, ensuring changes confined to Groovy pipeline code are recognized as non-therock-relevant and skip the TheRock CI pipeline. ## Test Plan No code logic was changed. Verified that `therock_configure_ci.py` pattern list is consistent with the existing `vars/*` skip entry and that the new pattern follows the same glob convention. ## Test Result N/A — directory scaffolding and CI filter only; no functional code affected. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation This pull request makes a small update to the `logging_mode_gtest.yaml` test configuration. The change narrows the GPU architecture pattern from `'12??'` to `'120?'`, making the test more specific in targeting GPU architectures. ROCM-25049
…new tests from test filter standardization (ROCm#8188) ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> New tests were added to hipcub, rocprim, and rocthrust without setting their RESOURCE_GROUPS properties. This may cause test failures when running ctest in parallel, e.g., `--parallel 8 --resource-spec-file resources.json`. This PR adds the missing RESOURCE_GROUPS properties. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> Without RESOURCE_GROUPS properties specified, ctest may launch these new tests all at once, causing them to compete on GPU resources. This often leads to failed/crashed/aborted tests and CI failures. ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> Run `ctest --parallel 8 --resource-spec-file resources.json` and all tests should pass. ## Test Result <!-- Briefly summarize test outcomes. --> All tests pass on local testing. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Song <hsong@ctr2-alola-ctrl-01.amd.com>
## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Add FFM test filters. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> This leverages the test filter standardization work. ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ctest --print-labels to see the new test categories/labels added ctest -L ffm-quick to run tests in the new category ## Test Result <!-- Briefly summarize test outcomes. --> Tests in the specified category are executed. Tests not in the specified category do something like "Running 0 tests from 0 test suites" - they are not executed. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Song <hsong@ctr2-alola-ctrl-01.amd.com>
## Motivation Some rocSPARSE pre_checkin tests were outside our time constraints significantly. This PR reduces the pre_checkin test times for the routines: bsrsm, bsrgeam, bsrilu0, spsm_csr, spsm_coo, spmm_csr, and gtsv
## Summary
Several CK source files carry Windows **CRLF** line endings (a trailing
carriage return on each line), introduced by editors configured for
Windows endings or copy/paste from Windows tooling. These are purely
cosmetic but they pollute diffs (whole-file churn the first time someone
makes an LF edit), confuse `clang-format`, and are inconsistent with the
LF-only convention used across the rest of the tree.
This PR (a) normalizes every existing CRLF file (6 files) to LF and (b)
adds a pre-checkin gate so new CRLF leaks are rejected before merge.
## File extensions covered
Both the cleanup scan and the new Jenkins enforcement stage use the same
predicate as the adjacent `ASCII Only Check` stage:
```
*.h *.hpp *.cpp *.h.in *.hpp.in *.cpp.in *.inc *.cl
```
(excluding `*/build/*` and `*/include/rapidjson/*`). The local
pre-commit hook's `c++/inc` type filter covers the same set.
## Why no enforcement today
CK is opted out of the rocm-libraries root `.pre-commit-config.yaml`, so
the existing `pre-commit` workflow doesn't touch CK. The local CK
`.pre-commit-config.yaml` only runs for developers who installed hooks.
The **authoritative gate is therefore the new Jenkins stage** in this
PR; the local hook is convenience.
## Commit layout (bisect-friendly)
1. `[ck] Normalize CRLF line endings to LF in C/C++ sources`
Mechanical line-ending cleanup across 6 files. No content change: every
edit is purely CRLF -> LF, verified with `git diff --ignore-cr-at-eol`
reporting an empty diff.
2. `[ck] Enforce LF-only line endings in C/C++ sources`
- New `projects/composablekernel/script/check_no_crlf.sh` (modeled on
`check_ascii_only.sh`).
- New `crlf-checker` entry in
`projects/composablekernel/.pre-commit-config.yaml` under the
local-hooks block (`types_or: [c++, inc]`).
- New `CRLF Check` parallel stage in
`projects/composablekernel/Jenkinsfile`'s `Static checks` block,
mirroring the adjacent `ASCII Only Check` stage. Always-on, no
`RUN_CPPCHECK` gate.
The tree is buildable at every commit boundary. Commit 1 leaves 0 CRLF
violations; commit 2 wires the gate.
## Demo
Script output on a synthesized violation:
```
$ printf 'int main() {}\r\n' > /tmp/bad.cpp
$ projects/composablekernel/script/check_no_crlf.sh /tmp/bad.cpp
ERROR: /tmp/bad.cpp contains CRLF (Windows) line endings:
1:int main() {}<CR>
Fix: convert to LF, e.g. 'sed -i 's/\r$//' /tmp/bad.cpp' or 'dos2unix /tmp/bad.cpp'
$ echo $?
1
```
Full repo scan after the cleanup commit:
```
$ cd projects/composablekernel && find . -type f \( -name '*.h' -o -name '*.hpp' -o -name '*.cpp' \
-o -name '*.h.in' -o -name '*.hpp.in' -o -name '*.cpp.in' -o -name '*.inc' -o -name '*.cl' \) \
-not -path '*/build/*' -not -path '*/include/rapidjson/*' -print0 \
| xargs -0 -P 8 -n 64 script/check_no_crlf.sh
$ echo $?
0
```
## Test plan
- [ ] Jenkins PR build: confirm new `Static checks -> CRLF Check` stage
runs green over the full predicate and the existing `ASCII Only Check` /
`Clang Format` stages are unaffected.
- [ ] Local: `pre-commit run crlf-checker --all-files` runs cleanly
after installing CK pre-commit hooks.
- [ ] Manually inject a CRLF line ending in any `.cpp/.hpp/.inc` file,
push: confirm Jenkins fails the new stage with a clear error.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
---------
Co-authored-by: Claude Opus 4 (1M context) <noreply@anthropic.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
## Motivation Look-back scan uses flags to communicate the state of blocks. Active polling is used to keep checking until a block either is in a partial state or a complete state. However, in the current implementation this is modeled as a per thread process. Meaning, each thread will keep polling global memory until only itself is in a valid state. But the entire wavefront only exits this loop when all threads have found a valid state. This can be potentially inefficient. While a block may present itself as partial, it may still be promoted to a complete block (but this will be invisible to the current implementation). And since the entire wavefront may still be working as long as an invalid state exists, it may unnecessarily look back onto other previous states. ## Technical Details The change here is to change flag getting to be a wave-cooperative and only exiting the active polling loop if all lanes have found a valid flag (via the `warp_any()` crosslane operator). Also the `get()` utility did duplicate flag retrieval, so the flag retrieval and other operations in the function are separated. ## Test Result <img width="1200" height="540" alt="device_scan_942" src="https://github.com/user-attachments/assets/6bd4695a-96b3-4a67-aa7d-ee8b03034187" /> <img width="1200" height="420" alt="device_scan_deterministic_942" src="https://github.com/user-attachments/assets/7af78c61-24d6-4463-a6b1-5d1e291b3fab" /> [lookback_scan.zip](https://github.com/user-attachments/files/26824239/lookback_scan.zip) ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Sander Bos <sander@streamhpc.com>
## Motivation Ragged tensors are tensors with variable per batch dimensions derived from an associated ragged offset tensor. This RFC outlines a plan for integrating them into the hipdnn frontend and flatbuffers, as well as presenting a design for utility classes which would provide users with the ability to allocate ragged tensors with proper indexing schemes. Furthermore, it details how these classes fit into the CPU reference validation integration testing framework. The immediate need for this work is expanding support for the SDPA forward CPU references so that they can match the functionality of AITER's SDPA forward kernels. With that in mind, this RFC uses those operations as an example on how this integration fits in with this process. ## Technical Details Technical details included in RFC ## Test Plan Docs only, no testing needed ## Test Result Docs only, no testing needed ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level None, this PR only adds documentation. **Associated ticket**: ALMIOPEN-1958
…m#8220) ## Motivation On the RetinaNet shapes (gfx950, fp16) CK Tile backward-data conv was ~18% behind classic CK, with the gap concentrated in the K=2376 3x3 detection-head family where bwd_data spends most of its time. The WAVELET GEMM pipeline already gives uplift for forward and backward-weight conv; this ports it to backward-data and consolidates the now-shared machinery across all three directions. ## Technical Details - Backward-data wavelet support in the tile kernel: launch extra load waves when the pipeline exposes `LaunchBlockSize`, and split the epilogue into math waves (run the CShuffle epilogue) and load waves (`RunBarrierStub`). - Register 7 WAVELET instances (fp16 and bf16), tuned for backward-data's tall-skinny GEMM rather than the forward tile shapes: a big-M `256/128/64` workhorse, a `VecA=4` variant for the `K % 8 != 0` shapes, and a `NumGroupsToMerge=32` variant for grouped (depthwise-style) shapes. - Implement the native backward-data instance parser in `generate_instances.py`. - Deduplicate the wavelet machinery shared by forward, backward-data, and backward-weight: `GroupedConvLaunchBlockSize`, `is_wavelet_pipeline`, and `RunWaveletAwareEpilogue` in `grouped_convolution_utils.hpp`; the three native instance parsers collapse to one parameterized parser. The three kernels now call the shared helpers. ## Test Plan - Rebuild the full profiler instance pools for all three directions (fp16/bf16/fp32, nhwgc/ndhwgc) to exercise the shared helpers across every instantiation. - Tile GTests on gfx950: `test_grouped_convnd_fwd_tile`, `test_grouped_convnd_bwd_data_tile`, `test_grouped_convnd_bwd_weight_tile`. - Per-shape sweep of the 35 RetinaNet backward-data shapes vs classic CK and the non-wavelet tile pool (`profile_wavelet_bwd_data.py`); correctness spot-checked with GPU-reference verification on the new big-M and NumGroupsToMerge instances. ## Test Result - GTests pass: forward 9/9, backward-data 6/6, backward-weight 6/6. - Backward-data perf (3x3 g=1 region, geomean classic/tile): 0.88 -> 1.11, i.e. the tile path goes from ~12% slower than classic to ~8% faster. The largest single backward-data shape (256x100x100->2376) moves from 11% slower than classic to 12.5% faster. - The dedup refactor preserves behavior (net -174 lines across the kernels/generator), confirmed by the full rebuild and the GTests above. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
## Motivation Dependency parser (dapper) requires the ninja build system. ## Technical Details This changes all CI builds over to ninja. Risk is low since it is very nearly a drop-in replacement for `make`. Build times are typically faster than `make`. ## Test Plan MICI. ## Test Result <TBD, should pass> ## Submission Checklist - [X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: SreecharanGundaboluAMD <sgundabo@amd.com>
## Motivation Fix a parameter type mismatch in TensileLite StreamK metadata handling so validation and serialization stay consistent with `validParameters`, eliminating avoidable YAML parameter type mismatch warnings in StreamK runs. ## Technical Details - Updated `tensilelite/Tensile/SolutionStructs/Solution.py` so `DirectToLdsMetadata` assignments use integer `0`/`1` instead of boolean `False`/`True`. - This aligns Solution state updates with the existing `validParameters` definition (`DirectToLdsMetadata: [0, 1]`) and avoids bool/int drift during Python/msgpack serialization. - Updated `tensilelite/Tensile/Tests/common/streamk/sk_dynamic.yaml` typing to match expected schemas for remaining mismatch fields (`MIArchVgpr: [False]`, `GlobalReadPerMfma: [1.0]`; commit `b7aeb1b83c`). - `sk_hybrid.yaml` was not validated on this branch because `StreamK=5` is unsupported on `develop`. ## Test Plan - Run unit validation test focused on parameter typing: - `tox -e unit -- Tensile/Tests/unit/test_validateParameterTypes.py -v` - Run StreamK dynamic client flow with `sk_dynamic.yaml` and inspect logs for YAML parameter type mismatch warnings. ## Test Result - `tox -e unit -- Tensile/Tests/unit/test_validateParameterTypes.py -v` -> `60 passed`. - `sk_dynamic.yaml` client run completed with exit code `0`. - No YAML parameter type mismatch warnings for `DirectToLdsMetadata`, `GlobalReadPerMfma`, or `MIArchVgpr` after this fix. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
## Motivation Port the `ConvOclDirectFwd` direct convolution solver from OpenCL to HIP, as part of the ongoing effort to remove the OpenCL backend dependency from MIOpen. ## Technical Details Add a new `ConvHipDirectFwd` solver and `MIOpenConvDirUniHip.cpp` HIP kernel that replaces the OpenCL-based `ConvOclDirectFwd` solver and `MIOpenConvDirUni.cl` kernel. The old solver and kernel are removed. Databases are updated by just replacing the solver name without updating the measurements because we need to run the tuning and DB updating round after all OCL-to-HIP conversions are completed. **New files:** - `src/kernels/MIOpenConvDirUniHip.cpp` — HIP kernel ported from the OCL `MIOpenConvDirUni.cl` kernel - `src/solver/conv/conv_hip_dir2Dfwd.cpp` — HIP solver ported from `conv_ocl_dir2Dfwd.cpp` - `src/solver/conv/conv_hip_dir2Dfwd_exhaustive_search.cpp` — exhaustive search ported from `conv_ocl_dir2Dfwd_exhaustive_search.cpp` **Removed files:** - `src/kernels/MIOpenConvDirUni.cl` — replaced by `MIOpenConvDirUniHip.cpp` - `src/solver/conv/conv_ocl_dir2Dfwd.cpp` — replaced by `conv_hip_dir2Dfwd.cpp` - `src/solver/conv/conv_ocl_dir2Dfwd_exhaustive_search.cpp` — replaced by `conv_hip_dir2Dfwd_exhaustive_search.cpp` **Modified files:** - `src/include/miopen/conv/solvers.hpp` — rename `ConvOclDirectFwd` → `ConvHipDirectFwd` and `ConvOclDirectFwdLegacyExhaustiveSearch` → `ConvHipDirectFwdLegacyExhaustiveSearch`; add backward-compat aliases for the fused solver - `src/solver.cpp` — register `ConvHipDirectFwd` with a new ID at end; replace old `ConvOclDirectFwd` registration with `++id` placeholder - `src/mlo_dir_conv.cpp` — replace `ConvOclDirectFwd` with `ConvHipDirectFwd` in the direct solver container - `src/fin/fin_interface.cpp` — update legacy ID 11 mapping to use `ConvHipDirectFwd` - `src/CMakeLists.txt` — add new source files, remove old ones - `docs/how-to/debug-log.rst`, `docs/reference/env_variables.rst` — update env var name `MIOPEN_DEBUG_CONV_DIRECT_OCL_FWD` → `MIOPEN_DEBUG_CONV_DIRECT_HIP_FWD` **Tests:** - `test/gtest/unit_conv_solver_ConvHipDirectFwd.cpp` — renamed from `unit_conv_solver_ConvOclDirectFwd.cpp`, updated solver references - `test/gtest/unit_FinInterface.cpp` — updated solver name and ID ## Test Plan - Ran unit tests verifying solver applicability and correctness (`unit_conv_solver_ConvHipDirectFwd`) - OCL-vs-HIP performance comparison benchmarks across ~14k shapes from FDB on gfx90a (MI-210) covering FP32, FP16, and BF16 — **is being collected** ## Test Result Benchmark results for gfx90a. We have a noticeable performance gain after OCL to HIP conversion for FP16. ### FP16 | Ratio (OCL/HIP) | Count | % | |------------------|------:|--:| | < 0.5 | 0 | 0.0% | | 0.5 - 0.7 | 36 | 0.6% | | 0.7 - 0.9 | 189 | 3.3% | | 0.9 - 1.0 | 725 | 12.8% | | 1.0 - 1.1 | 707 | 12.5% | | 1.1 - 1.3 | 991 | 17.6% | | 1.3 - 1.5 | 835 | 14.8% | | 1.5 - 2.0 | 559 | 9.9% | | 2.0 - 3.0 | 967 | 17.1% | | > 3.0 | 636 | 11.3% | | Metric | Value | |--------|-------| | Total shapes | 5645 | | Mean ratio | 1.6784 | | Min ratio | 0.5051 | | Max ratio | 3.3972 | | HIP regression (< 0.9) | 225 (4.0%) | | Parity (0.9 - 1.1) | 1432 (25.4%) | | HIP faster (> 1.1) | 3988 (70.6%) | | Shapes with nonzero max_abs_err | 0 (0.0%) | ### BP16 - to be added | Ratio (OCL/HIP) | Count | % | |------------------|------:|--:| | < 0.5 | 185 | 4.7% | | 0.5 - 0.7 | 65 | 1.6% | | 0.7 - 0.9 | 686 | 17.3% | | 0.9 - 1.0 | 1242 | 31.3% | | 1.0 - 1.1 | 832 | 21.0% | | 1.1 - 1.3 | 487 | 12.3% | | 1.3 - 1.5 | 345 | 8.7% | | 1.5 - 2.0 | 78 | 2.0% | | 2.0 - 3.0 | 23 | 0.6% | | > 3.0 | 24 | 0.6% | | Metric | Value | |--------|-------| | Total shapes | 3967 | | Mean ratio | 1.0279 | | Min ratio | 0.2912 | | Max ratio | 7.8593 | | HIP regression (< 0.9) | 936 (23.6%) | | Parity (0.9 - 1.1) | 2076 (52.3%) | | HIP faster (> 1.1) | 955 (24.1%) | | Shapes with nonzero max_abs_err | 0 (0.0%) | ### FP32 - to be added ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com> Co-authored-by: BradPepersAMD <Brad.Pepers@amd.com> Co-authored-by: Daming Feng <dmfeng8898@gmail.com>
## Summary This PR relaxes BatchNorm frontend validation to allow broadcastable shapes for scale and bias tensors, addressing issue ROCm#6409. Previously, BatchNorm required scale and bias to strictly follow the shape `{1, C, 1, 1}`, enforced via `validateChannelOnlyTensorShape`. This restriction is inconsistent with other normalization operators such as LayerNorm and the recently updated RMSNorm (PR ROCm#6359), both of which allow more flexible, broadcastable shapes. This change aligns BatchNorm with the same design philosophy, making the frontend more permissive while leaving backend-specific constraints to the plugin layer. --- ## Changes ### Validation Updates - Removed strict `{1, C, 1, 1}` requirement for scale and bias tensors - Replaced channel-only validation with: - Minimum dimension validation (`>= 1`) - Shape consistency check between scale and bias - Now supports broadcastable shapes such as: - `{C}` - `{1, C}` - `{1, C, 1, 1}` (previously required format) ### Preserved Constraints - Scale and bias must still have **matching shapes** - Mean and inverse variance tensors remain **strictly channel-shaped** - Input/output tensor shape consistency is unchanged - Spatial validation (`N * spatial > 1`) remains intact - Running statistics validation logic is unchanged --- ## Tests - Added tests to cover broadcastable scale/bias shapes: - Rank-1 (`{C}`) - Rank-2 (`{1, C}`) - Ensured existing tests reflect relaxed frontend validation - Removed outdated test enforcing strict channel-dimension matching - Verified all edge cases: - Shape mismatches - Missing attributes - Spatial constraints - Running statistics - 4D and 5D inputs All tests pass locally (22/22) --- ## Design Considerations - Frontend validation is intentionally permissive: - Accepts broadcastable shapes - Does **not** enforce backend-specific constraints - Backend providers (e.g. MIOpen) remain responsible for: - Validating supported tensor formats - Rejecting unsupported broadcast cases This separation ensures: - Consistency across normalization operators - Flexibility for future backend improvements --- ## Notes - MIOpen currently supports only `{1, C, 1, 1}` scale/bias tensors - Broader shape support will require future updates in the miopen-provider and MIOpen itself --------- Co-authored-by: BrianHarrisonAMD <169072757+BrianHarrisonAMD@users.noreply.github.com>
…OCm#8262) ## Motivation Add HIP graph capture support for FMHA backward operations. The original implementation only supported normal execution mode and would cause use-after-free crashes when used with graph capture replay. When FMHA backward is captured into a HIP graph: - First replay: host callback executes and deletes the closure (as designed for normal mode) - Subsequent replays: use-after-free crash because the closure was already freed This PR enables `fmha_bwd_launcher::prepare_workspace_async()` to work correctly in both normal execution and graph capture modes.
…neLibrary and TensileUpdateLibrary (ROCm#7703) ## Motivation Continuation of the TensileLite Python unit-test coverage uplift: this PR adds pytest suites for four library-management modules (`TensileLogic/Run`, `TensileMergeLibrary`, `TensileRetuneLibrary`, `TensileUpdateLibrary`), all at 0% coverage on `develop`. **In addition**, this PR makes three product-code changes in `TensileRetuneLibrary.py` and `TensileUpdateLibrary.py` that were either prerequisites for unit-testing the modules in isolation (the import refactor) or bugs/dead code uncovered while writing the tests. All three are detailed below. Tracker (internal): https://amd-hub.atlassian.net/browse/AIHPBLAS-3405. Pairs with ROCm#7722 — see "Cross-PR coupling" below. ## Product code changes This PR is **not** tests-only. Three product-code changes are bundled: 1. **Granular import refactor** (`TensileRetuneLibrary.py`, `TensileUpdateLibrary.py`). Replaced the umbrella `from .Common import …` lines with specific submodule imports — `from .Common.GlobalParameters import …`, `from .Common.Utilities import …`, `from .Common.Constants import …`, etc. **Rationale:** lets the modules be unit-tested without pulling the full `Common` package and its hardware-side initialization. **Safety:** every name removed from the umbrella is either reachable via the new granular imports or still re-exported from `Common/__init__.py` via existing `from .Constants/.Utilities/.Parallel import *` lines (and ROCm#7722 explicitly re-exports the four `GlobalParameters` names — see Cross-PR coupling). 2. **Dead-code removal** (`TensileRetuneLibrary.py`). Removed an exact-duplicate copy of `pushWorkingPath`, `popWorkingPath`, `ensurePath`, and `setWorkingPath` (defined twice in the file on `develop`, lines 45–66 and again at 69–90 with identical bodies — the second copy was unreachable shadowing). **Safety:** the surviving copies are byte-for-byte identical to the removed ones. 3. **Bug fix** (`TensileRetuneLibrary.py::parseCurrentLibrary`). `GlobalParameters.globalParameters["PerformanceMetric"] = libYaml[10]` → `globalParameters["PerformanceMetric"] = libYaml[10]`. The qualified form referenced a name (`GlobalParameters`) that is never imported in this module — it would have raised `NameError` if the optional 11th element of `libYaml` ever appeared in a real config. **Locked by test:** `Tests/unit/test_TensileRetuneLibrary.py::TestParseCurrentLibrary::test_parses_library_without_size_file` mocks `LibraryIO.read` to return an 11-element list, forcing the `if len(libYaml) > 10:` branch. ## Cross-PR coupling This PR pairs with **ROCm#7722**, which adds `from .GlobalParameters import globalParameters, assignGlobalParameters, restoreDefaultGlobalParameters, __version__` to `Tensile/Common/__init__.py`. That re-export preserves the historical `from Tensile.Common import globalParameters` path for any other internal or external caller after this PR's refactor lands. The two PRs are reviewable independently but ideally land together; if landed in either order, no caller is broken (existing `*`-imports in `Common/__init__.py` already cover the other names). ## Technical Details **Modules covered in this PR:** | Module | Before | After | |--------|-------:|------:| | `Tensile/TensileLogic/Run.py` | 0.00% | **94.86%** | | `Tensile/TensileMergeLibrary.py` | 0.00% | **96.40%** | | `Tensile/TensileRetuneLibrary.py` | 0.00% | **94.87%** | | `Tensile/TensileUpdateLibrary.py` | 0.00% | **95.65%** | Project-wide TensileLite total: 24.66% → **25.68%**. **Testing approach:** - Pure-Python unit tests, no code-generator invocation. - Extensive mocking with `unittest.mock` for all external dependencies (file I/O, YAML read/write, `validateToolchain`, `ClientWriter`, `LibraryLogic`, `ProblemType`, `Solution`, `ProblemSizes`). - All tests marked `@pytest.mark.unit`. - Coverage targeted at normal execution paths, error conditions, edge cases, state changes, and CLI-argument parsing. Reference commit for the test additions: ROCm@a62795d ## Test Plan Pure-Python unit tests, CPU only. Run via the existing tensilelite tox `unit` env: ```bash cd projects/hipblaslt/tensilelite tox -e unit -- \ Tensile/Tests/unit/test_TensileLogic_Run.py \ Tensile/Tests/unit/test_TensileMergeLibrary.py \ Tensile/Tests/unit/test_TensileRetuneLibrary.py \ Tensile/Tests/unit/test_TensileUpdateLibrary.py ``` The bug fix in `parseCurrentLibrary` is exercised specifically by `test_TensileRetuneLibrary.py::TestParseCurrentLibrary::test_parses_library_without_size_file`. ## Test Result Tests pass; new modules at the coverage levels listed in Technical Details. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Signed-off-by: pdhirajkumarprasad <dhirajp@amd.com>
## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> Add FFM test filters. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> This leverages the test filter standardization work. ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> ctest --print-labels to see the new test categories/labels added ctest -L ffm-quick to run tests in the new category ## Test Result <!-- Briefly summarize test outcomes. --> Tests in the specified category are executed. Tests not in the specified category do something like "Running 0 tests from 0 test suites" - they are not executed. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Song <hsong@ctr2-alola-ctrl-01.amd.com>
…#8204) ## Motivation <!-- Explain the purpose of this PR and the goals it aims to achieve. --> SIA0+PGR2+PLR0 TDM kernels gave wrong results because the cross-wave TDM buffer swap barrier didn't wait for in-flight local reads to drain ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> Insert a dscnt=0 wait before the TDM swap _syncThreads barrier when numItersPLR==0, ensuring all local reads finish before the buffer is reused. ## Test Plan <!-- Explain any relevant testing done to verify this PR. --> Add a case ## Test Result <!-- Briefly summarize test outcomes. --> Pass tox test (those failed case are known issues: sk no solution) <img width="2557" height="215" alt="image" src="https://github.com/user-attachments/assets/1e0fb8b2-2513-4c7a-844c-dba6fb97898f" /> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
Complex acc VGPRs interleave real/imag, so the relative offset elementSumIdx[i]-elementSumIdx[0] cannot locate the imag half. Disable skipRearrangement for complex to use the correct reorder path.
… SK5 (ROCm#8568) ## Motivation Switch gfx950 device library logic YAML files (excluding MX kernels) from SK3 to SK5 with default-OFF mode. SK5-OFF executes the SK3 code path at runtime but enables future hybrid-mode scheduling when turned ON. **Depends on:** PR ROCm#8162 (SK5 kernel infrastructure in TensileLite) ## Changes - gfx950 Equality, Origami, GridBased, and Range YAML files (519 files) updated: - `StreamK: 3` → `StreamK: 5` in solution parameters - `_SK3_` → `_SK5_` in kernel/solution name strings - No host-side code changes - No custom kernel assembly changes **MX kernels are NOT changed** — both F4/MXFP4 microscaling (`F4*_MXA32_MXB32`) and matrix-bias (`S_MX_B`) variants remain on SK3, as they do not yet support the SK4 path used by SK5 hybrid mode. ## Change Breakdown by Directory | Directory | Files | ~Lines Changed | Description | |-----------|-------|----------------|-------------| | Equality (gfx950) | 23 | ~46K | Core equality logic (excl. MX) | | Equality (gfx950_id75a3) | 23 | ~38K | id75a3 variant (excl. MX) | | Equality (gfx950_id75a8) | 3 | ~4K | id75a8 variant | | Origami (root) | 153 | ~113K | Origami base (excl. MX) | | Origami_nta4 | 153 | ~113K | Origami NTA4 (excl. MX) | | Origami_ntb4 | 153 | ~113K | Origami NTB4 (excl. MX) | | GridBased | 8 | ~2K | Grid-based solutions | | Range | 3 | 36 | Range-based solutions | | **Total** | **519** | **~430K** | | ## Split PRs for Review Since GitHub cannot render diffs for 519 files / 430K lines, the PR is split into 4 reviewable parts: | Part | Content | Files | PR | |------|---------|-------|----| | 1 | Equality + GridBased + Range | 60 | ROCm#8596 | | 2 | Origami base | 153 | ROCm#8597 | | 3 | Origami_nta4 | 153 | ROCm#8598 | | 4 | Origami_ntb4 | 153 | ROCm#8599 | All YAML changes are the same mechanical substitution (`StreamK: 3` → `5`, `_SK3_` → `_SK5_`). Reviewing Part 1 in detail is sufficient; Parts 2-4 are the same pattern applied to Origami files. ## Branch Verification The 4 split branches are merged into a verification branch. If the compare below shows **no diff**, the split is complete and correct: **[Compare: main branch vs merge of all 4 parts](ROCm/rocm-libraries@users/jolabega/sk5-device-library-default-off...users/jolabega/sk5-default-off-merge-verify)** ```mermaid graph LR develop[origin/develop] --> part1[Part 1: Equality] develop --> part2[Part 2: Origami root] develop --> part3[Part 3: Origami_nta4] develop --> part4[Part 4: Origami_ntb4] part1 --> merged[merge-verify branch] part2 --> merged part3 --> merged part4 --> merged merged -->|"compare shows no diff"| main[Main PR branch] ``` ## Test Plan - CI device library validation (`TensileLogic --check-all`) - `hipblaslt-test` pre_checkin suite on gfx950 ## Submission Checklist - [X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level Low — changes are limited to device library logic YAML files (kernel selection metadata). No host code, no algorithm changes, no assembly changes. MX kernels excluded. **Associated ticket:** AIHPBLAS-3921
…n data (ROCm#8522) ## Summary The SDPA forward kernel supports an optional LSE (Log-Sum-Exp) statistics output tensor, but the golden data generator ignores the `--stats` flag and the C++ plan builder rejects any graph that requests it. This blocks backward golden data generation because the backward kernel requires LSE as an input tensor. This PR wires up LSE output end-to-end across the Python generator, provider C++, CPU reference executor, and integration tests. ## Risk Assessment Low risk. Changes are scoped to test infrastructure (golden data generator, CPU reference executor, test harness) and provider forward plan internals behind an existing `generate_stats` gate. No public API changes. Unit and integration tests pass with no regressions. ## ASIC Coverage Standard PR CI is sufficient. The changes affect test infrastructure (golden data generation, CPU reference execution) and provider-internal plan building behind an existing feature gate. No kernel selection, support surface, or default behavior changes. ## Testing Summary - Unit tests: 305/305 hip-kernel-provider tests pass with no regressions and backward compatibility verified. - Integration tests: 12/12 CPU golden reference tests pass across quick and standard tiers for BF16/FP16 with and without stats. - Cross-validation: AITER v3 forward cross-validation passed for all 6 BF16 bundles (max_abs < 0.001). - Backward compatibility: `--stats` omitted produces unchanged output (4 tensors, `generate_stats: null`). ## Testing Checklist - [x] hip-kernel-provider unit tests - `./build/bin/hip_kernel_provider_tests` - Status: Passed - [x] CPU golden reference integration tests - `hipdnn_integration_tests --gtest_filter="*CpuSdpaFwd*"` - Status: Passed - [x] AITER v3 cross-validation - 6 BF16 bundles, max_abs < 0.001 - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - Add `compute_lse()` to Python generator to produce reference LSE tensor (uid=4, shape `[B, H_q, S_q, 1]`, FP32) via `torch.logsumexp`, wired to `--stats` flag. - Fix FP16 dtype string (`"half"` not `"float16"`) and AITER version detection in generator; bump version to 1.0.1. - Add `lseUid` and `lseStrideHead` to `SdpaFwdParams` and wire conditional LSE pointer/stride logic in `SdpaFwdPlanBuilder` and `SdpaFwdPlan`. - Update CPU reference executor to accept and compute rank-4 LSE `[B, H, Sq, 1]` instead of rejecting stats graphs. - Introduce `PlanNotApplicableException` typed exception for structured error handling when a CPU plan builder is not applicable. - Reorganize golden data bundles into `quick/` and `standard/` tiers with FP16 and stats variants (12 bundles total, DVC-pushed). --------- Co-authored-by: Claude Opus 4 <noreply@anthropic.com>
…ucket; fix Linux bindings preload (ROCm#8581) ## Summary Makes the dnn-benchmarking ROCm PyTorch setup work on gfx90a (MI200/MI210/MI250) and makes the bare-image build self-contained from a single input. Three changes: 1. **Nightly bucket** — `setup.sh` selected the `gfx90X-dcgpu` bucket, frozen at torch 2.12 / ROCm 7.12 (predates several ROCm SDK libraries, notably hipDNN). Point gfx90a at the current `gfx90a` bucket (torch 2.13 / ROCm 7.14, which ships hipDNN). 2. **`--gpu-arch` as single source of truth** — `setup.sh` now derives the HIP offload target from `--gpu-arch` and passes `-DGPU_TARGETS` to the hipDNN/provider cmake builds, and its torch-mode probe tolerates the GPU-less `import torch` SDK warning. A bare/GPU-less build (e.g. `docker build`) no longer needs external `GPU_TARGETS`/`AMDGPU_TARGETS`/`ROCM_SDK_TARGET_FAMILY` env or GPU-discovery shims. 3. **Linux bindings preload** — preload the base HIP runtime via `rocm_sdk` instead of the `hipdnn` shortname. ## Risk Assessment Low risk. Changes are confined to the dnn-benchmarking `setup.sh` and the hipDNN Python bindings' import-time preload list. gfx942/gfx950 bucket selection is unchanged. The new arch handling uses hipDNN's documented `-DGPU_TARGETS` interface and only affects the device-code builds (the nanobind bindings are host-only). All paths are verified end-to-end on MI210 in a clean Ubuntu 24.04 + Python 3.12 environment built with no ROCm-related env beyond `--gpu-arch`. The Windows preload list is unchanged (still names `hipdnn`) but was not exercised. ## Testing Summary - Full dnn-benchmarking suite from a clean Ubuntu 24.04 + Python 3.12 install following the documented manual steps, passing only `--gpu-arch` (no `GPU_TARGETS`/`ROCM_SDK_TARGET_FAMILY` env, no offload shims). Validates that setup.sh's `-DGPU_TARGETS` device codegen builds and runs on the MI210. - GPU/ROCm-marked subset, exercising hipDNN-backed execution through the engine plugins. - Import smoke with venv activation only (no `LD_LIBRARY_PATH` workarounds). - `get_torch_mode` final-line parsing unit-checked against warning-polluted stdout (incl. a warning with no trailing newline) and against a plain mode-only interpreter. ## Testing Checklist - [x] dnn-benchmarking suite (bare Ubuntu 24.04, `--gpu-arch` only) - `pytest -m "not cuda"` - ASICs: gfx90a - Status: Passed (815 passed, 8 skipped, 8 xfailed) - [x] GPU/ROCm subset - `pytest -m "gpu or rocm"` - ASICs: gfx90a - Status: Passed (55 passed) - [x] setup.sh unit tests (incl. existing-CUDA-venv rejection) - `pytest tests/unit/cli/test_setup_script.py` - Status: Passed (4 passed) - [x] Bindings import, venv-only - `python -c "import hipdnn_frontend"` - ASICs: gfx90a - Status: Passed - [ ] Windows ROCm-wheel import - Status: Not run - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - `setup.sh` (bucket): map `gfx90a` to the `gfx90a` ROCm nightly bucket (current; ships hipDNN) instead of the frozen `gfx90X-dcgpu` family bucket. `gfx942`/`gfx950` keep their `-dcgpu` buckets. - `setup.sh` (arch ownership): resolve the GPU arch once (`--gpu-arch` or detection) and pass `-DGPU_TARGETS`/`-DAMDGPU_TARGETS` to the `build_hipdnn`/`build_provider` cmake invocations; export `PYTORCH_ROCM_ARCH` as belt-and-suspenders. The wheel SDK ships no `rocm_agent_enumerator`/`offload-arch` on PATH and the build may run with no GPU, so HIP cannot autodetect the offload arch — supplying it explicitly removes the need for caller-set env and shims. - `setup.sh` (torch-mode): `get_torch_mode()` prints the mode on its own final line and the caller reads only that last line, so the GPU-less `import torch` SDK warning no longer corrupts torch-mode detection (a plain mode-only interpreter still works). Removes the `ROCM_SDK_TARGET_FAMILY` workaround. - `hipdnn_frontend/__init__.py`: on Linux, preload the base HIP runtime (`amd_comgr`, `amdhip64`, `hiprtc`) instead of `hipdnn` (`rocm_sdk` resolves base ROCm SDK packages; the libraries prefix is already on `LD_LIBRARY_PATH`). Windows keeps the full list including `hipdnn`.
## Motivation Optimize pipeline V3 for gfx950 by enabling buffer load to lds (async pipeline) ## Technical Details - Add `Async` bool to `Problem` struct to enable async pipeline in existing one - Add `static_move_ys` to load transpose. This generates offset in assembly instructions saving registers - Add `is_valid` to `async_get_vectorized_elements`. Before hard coded to true. It allows to support padding - Remove unnecessary restrictions to `is_a_load_tr` and `is_b_load_tr` (wider use of lds load transpose on gfx950) - Integrate async support in existing V3 pipeline (avoid pipelines duplication) - Create policy to support both async and default cases. This could be used by any async pipeline (next steps) - Define `wg_attr_num_access` separately for A and B. This allows to optimize ds_read instruction width for cases when one matrix is transposed and the other is not. Before in such cases, `ds_read_b64` was used instead of `ds_read_b128` - Add test for V3 async. Currently only supporting cases with A and B having the same type ## Test Plan New test `test_ck_tile_gemm_pipeline_compv3_async` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
…CPU probe + setup.ps1) (ROCm#8255) ## Summary Make the dnn-benchmark tool and its non-GPU test suite run on Windows. This replaces the Unix-only CPU-time probe with a cross-platform one, adds a Windows PowerShell setup script (`setup.ps1`), and updates platform-specific unit-test assumptions so the non-GPU suite passes on Windows. ## Risk Assessment Low risk. Changes are confined to the dnn-benchmark tool: a cross-platform rework of the CPU-time probe, an additive Windows setup script, and platform-aware test assertions. No public API, schema, or build-behavior changes; the non-GPU suite passes locally on Windows. ## Testing Summary - dnn-benchmark non-GPU unit/integration suite on Windows (CPython 3.12, CPU torch). - CLI entry-point smoke check. - `setup.ps1` run end-to-end into the Windows ROCm-wheel env. ## Testing Checklist - [x] dnn-benchmark non-GPU tests - `pytest -m "not gpu"` - Status: Passed (684 passed, 10 skipped) - [x] CLI smoke - `python -m dnn_benchmarking --help` - Status: Passed - [x] Windows setup script - `setup.ps1` against the ROCm-wheel env - Status: Passed - [ ] GPU tests - require the MIOpen provider engine plugins (Linux-only) - Status: Not run - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - `metrics/host.py`: sample process CPU time via `os.times` — the cross-platform stdlib accessor (backed by `GetProcessTimes` on Windows) — so `CpuTimeProbe` works on Windows and the module no longer depends on the Unix-only `resource` module. `_process_cpu_times()` returns a plain `(user, kernel)` tuple that the probe consumes positionally. The probe wraps the whole benchmark loop and reports a per-iteration average, so `os.times`'s clock-tick resolution is immaterial; it degrades to `None` if `os.times` is unavailable. - `setup.ps1` (new): a Windows PowerShell analogue of `setup.sh`. Installs dnn-benchmark into a selected Python env, optionally builds hipDNN + Python bindings + the MIOpen provider from source (`-ForceBuild`) inside an MSVC vcvars64 + Windows SDK environment, wires the compiled bindings via a `.pth`, installs PyTorch per `-TorchMode` (`cpu`/`existing`/`none`), and verifies the result. setup.sh's venv management (`--reuse-venv`/`--workspace`) and ROCm/CUDA torch modes are intentionally omitted (not applicable on Windows). - Tests: make platform-specific unit tests Windows-aware — compare paths via `Path` / `as_posix()` / `str(Path(...))` rather than hard-coded POSIX separators, embed a TOML fixture path as POSIX, update `test_host.py` for the tuple-based CPU-time sample, and skip two inherently Unix-only tests (a `?`-in-filename sqlite URI case and a `+x` executable-bit case that `os.access(X_OK)` cannot model on Windows). --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
## Motivation Switch to dispatcher profiler for ck tile conv. ## Technical Details - Switch to dispatcher profiler for ck tile conv. - Drop profiler for experimental codegen - Minor fixes for bwd data printing - Minor fixes for 3d conv in dispatcher codegen ## Test Plan test_grouped_conv*tile ## Test Result Passed ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Ville Pietilä <> Co-authored-by: Ville Pietilä <188998872+vpietila-amd@users.noreply.github.com>
) ## Motivation The Batchnorm backward activation integration test for the mlops engine in the hip-kernel-provider was only testing the NCHW layout. However, the engine also supports NHWC, NCDHW, and NDWC layouts for this operation which are tested in the other batchnorm integration tests for the engine. ## Technical Details This patch expands the testing to cover this gap. ## Test Plan New tests are added to ` hip_kernel_provider_integration_tests` ## Test Result Test pass on MI210 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
…OCm#8626) ## Motivation Update subtile heuristic to only restrict based on K<512.
[CI] Run hipSPARSELt when hipBLASLt subtree changes (ROCm#7514) Map projects/hipblaslt to also activate the sparselt optional matrix project so hipSPARSELt builds and tests in the blas TheRock job. Stop mutating module-level state in collect_projects_to_run (per-call deep copies) and update therock_matrix tests accordingly. Tracking: ROCM-25320 Fixes ROCm#7519 Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Yanyao Wang <yanywang@amd.com>
## Summary
Enable the first multi-threaded / multi-stream nightly test for
`hipblaslt-test`. Until now every test ran single-threaded (`threads:0,
streams:0`), which left the concurrent-usage paths that real customers
hit (JAX/XLA, MaxText, etc.) uncovered in CI. This PR turns on `{
threads: 3, streams: 3 }` in the shared `common_threads_streams` anchor
and adds one nightly matmul entry (`matmul_multithread_stream`) that
consumes it on gfx942 / gfx950.
This is purely a test-data change. The test harness already exposes
`threads` / `streams` arguments via `RUN_TEST_ON_THREADS_STREAMS` in
`clients/common/include/hipblaslt_test.hpp`; the underlying concurrency
correctness fixes landed in ROCm#5597 (`hipStreamSynchronize`,
`hipMemcpyAsync`, mutex on the `memory_pool` singleton).
## Work tracking
- Task (this PR): https://amd-hub.atlassian.net/browse/AIHPBLAS-1023
- Parent Epic: https://amd-hub.atlassian.net/browse/AIHPBLAS-1003 —
"hipBLASLt Testing & Quality Improvements — Prevent Production Crashes"
- Companion harness fix (merged): ROCm#5597 (AIHPBLAS-1431)
- Plan / failure log:
https://amd.atlassian.net/wiki/spaces/MLSE/pages/1524793988
(AMD-internal)
- Customer-visible motivators (AMD-internal): SWDEV-576540 (JAX/XLA
matmul segfault), SWDEV-565755 (JAX training docker blocker), AIXLA-208
(JAX free() segfault in ROCm 7.2)
## Tests
**Paths**
- `projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml` —
uncomments `{ threads: 3, streams: 3 }` in the `&common_threads_streams`
anchor. No existing test on `develop` references this anchor, so the
only behavioral effect is on the new entry below.
- `projects/hipblaslt/clients/tests/data/matmul_gtest.yaml` — adds
`matmul_multithread_stream` (nightly, gfx942/gfx950), mirroring
`matmul_medium` with `threads_streams: *common_threads_streams` added.
**Run**
```
./hipblaslt-test --gtest_filter='*matmul_multithread_stream*'
```
**Results (post-ROCm#5597, threads:3 / streams:3)**
- Targeted suite: 384 tests, ~32s, all PASSED.
- Full nightly (`*matmul_test*`) on gfx942: 22,435 tests, ~51 min, all
PASSED.
**Why `threads:3, streams:3`**
Matches the rocBLAS coverage convention. The `{ threads: 4, streams: 4
}` and `{ threads: 5, streams: 5 }` rows are intentionally left
commented because they currently surface unresolved failures
(`matmul_bad_arg`, several `matmul_heuristic_all_solutions` variants).
Those failures need separate tracking before they can be enabled — see
PR comment.
## Flags / guardrails
- [x] N/A — no product code changed; new coverage is nightly only and
limited to `gpu_arch: '9(42|50)'`.
## Adjacent tests considered
- Other operations (groupedgemm, ext-API, smoke, rocroller) —
deliberately out of scope for this first enable; can land incrementally
under AIHPBLAS-1023.
- Multi-device (`devices > 1`) — out of scope here; depends on multi-GPU
CI capacity called out in AIHPBLAS-1003.
- Higher concurrency (`threads/streams: 4` and `5`) — deferred behind
known failures, see Tests section.
## Risk acceptance
- [x] STANDARD — test-only, nightly category, gfx942/gfx950 only, no
flag needed.
## Submission Checklist
- [x] Looked over contributing guidelines
(https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests).
---------
Signed-off-by: pdhirajkumarprasad <dhirajp@amd.com>
Co-authored-by: Tony Davis <tony.davis@amd.com>
Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com>
## Motivation Selecting the tile size for attention algorithms. ## Technical Details Origami model for FA2. ## Test Plan Added additional tests for the attention model. ## Test Result All tests are passing ## Selection accuracy: 66.6% <= 8K problem sizes and 100% > 8K problem sizes ``` q_seq_len , kv_seq_len , head_dim , q_heads , kv_heads , batch , dtype , selected_block_m , selected_block_n 32 , 32 , 128 , 32 , 32 , 1 , bf16 , 32 , 32 64 , 64 , 128 , 32 , 32 , 1 , bf16 , 32 , 64 128 , 128 , 128 , 32 , 32 , 1 , bf16 , 64 , 64 256 , 256 , 128 , 32 , 32 , 1 , bf16 , 128 , 64 512 , 512 , 128 , 32 , 32 , 1 , bf16 , 128 , 64 1024 , 1024 , 128 , 32 , 32 , 1 , bf16 , 128 , 64 2048 , 2048 , 128 , 32 , 32 , 1 , bf16 , 128 , 64 4096 , 4096 , 128 , 32 , 32 , 1 , bf16 , 128 , 64 8192 , 8192 , 128 , 32 , 32 , 1 , bf16 , 128 , 64 ``` ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Ryan Swann <109695074+ryanswann-amd@users.noreply.github.com>
rocRoller was tested on every BLAS-group change (hipBLASLt, rocBLAS, etc.) because shared/rocroller mapped to the blas group and rocroller was in its projects_to_test. hipBLASLt no longer relies on rocRoller kernels, so this only added flaky/timeout-prone CI with no signal. Route shared/rocroller to its own rocroller group via additional_options with project_to_add="blas": rocRoller is now tested only when its own subtree changes, still builds under the BLAS umbrella, and merges into the blas job when a PR touches both. Also drop shared/rocroller from the "project: hipblaslt" labeler entry to match the new separation. Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Yanyao Wang <yanywang@amd.com>
ROCm#8205) ## Motivation The goal of this PR is to standardize the precision support reference page format across all components, while also reducing the maintenance of burden of having to manually update the YAML data file in https://rocm.docs.amd.com/en/latest/reference/precision-support.html ## Technical Details - Each component maintains its own YAML file which will be eventually used in https://rocm.docs.amd.com/en/latest/reference/precision-support.html - A new precision support reference page is introduced which will not override existing data type/precision support content; it will serve as the overview/summary that will be linked in the ROCm reference page ## Test Plan - Built locally, viewed each component manually ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation The existing integration test harness hard-coded a single verification path (CPU reference to golden comparison) with the load/compare logic scattered across `SetUp()`. This made it difficult to add new verification modes (GPU reference, engine-under-test) or onboard non-golden bundle formats. ALMIOPEN-1968 restructures the harness into a generic runner that separates discovery, loading, and verification so all three executor modes (CPU ref, GPU ref, engine) can share the same bundle infrastructure. ## Key changes - **Bundle discovery** (`BundleDiscovery.hpp`): walks an arbitrary directory tree, supports both the tiered golden layout and flat customer drops under one root - **`IntegrationTestBundle`** struct + `loadIntegrationTestBundle()`: consolidates the three-step pattern (`loadBundleMetadata` + `loadGraphAndTensors` + `extractAndClearOutputTensorData`) into a single call - **Pre-load checks** (`graphJsonParses`, `tensorDataPresent`): disambiguate malformed bundles (FAIL) from missing DVC data (SKIP), since `loadGraphAndTensors()` throws for both - **Rich failure reports**: tensor UID, shape, dtype, max abs/rel error vs tolerance, worst-element index with expected/actual values, and mismatch count - **Golden to generic rename**: `GoldenBundleDiscovery.hpp` to `BundleDiscovery.hpp`, `GoldenBundleRegistration.hpp` to `BundleRegistration.hpp`, `GoldenBundleLoadCheck.hpp` to `BundleLoadCheck.hpp` and all corresponding symbols — the infrastructure is bundle-generic, not golden-specific - **Metadata guards** (`applyMetadataGuards`): VRAM requirement and arch compatibility checks gate expensive tensor loads with early SKIP ## Testing - All 189 unit tests pass (169 green, 20 skipped for no GPU device) - `loadIntegrationTestBundle()` has direct unit tests: `LoadIntegrationTestBundlePopulatesAllFields` and `LoadIntegrationTestBundleThrowsOnMissingBin` - Discovery, load-check, and verification path tests all updated and passing --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
## Motivation Fix int8 GEMMs crashing with INTERNAL_ERROR when alpha is 1065353216 (the int32 bit pattern of float 1.0), and remove duplicated alpha/beta classification logic that let two code paths disagree. ## Technical Details Promote get_scalar_value_from_void_ptr to a shared helper so ConstructTensileProblem and updateTensileProblem both classify alpha/beta from the true storage type (alphaBetaType) instead of the matrix type / a magnitude double; also delete the dead get_alpha_beta_target_type. ## Test Plan Added regression test matmul_gemm_i8_dst_i32_alpha_float_bits (alpha = int32 bits of float 1.0 = 1065353216) on gfx942/gfx950; ran the complex and i8_dst_i32 suites. ## Test Result New test passes (4 cases); complex (2314) and i8_dst_i32 (188) suites pass — kernel selection succeeds for the int32 alpha that previously crashed. --------- Co-authored-by: Naveen Kumar Elumalai <nelumala@ctr2-alola-ctrl-01.amd.com>
…k-sparsity-merge # Conflicts: # projects/composablekernel/example/ck_tile/01_fmha/example_fmha_fwd.cpp # projects/composablekernel/example/ck_tile/01_fmha/fmha_fwd_runner.hpp # projects/composablekernel/test/ck_tile/fmha/test_fmha_fwd.cpp
goldcoderZ
pushed a commit
that referenced
this pull request
Jun 26, 2026
## Motivation
The current top of tree ctests for sanity testing (test 1 & test 2) do
not run from external directories. This PR will allow these tests to run
from any directory
## Technical Details
Use ctests built-in `build-and-test` instead of building separately
## Test Plan
Run ctests
## Test Result
From a directory outside of actual build directory
```
lakshmi@kapu:~/work/rpp-test$ ctest -VV -I 1,2,1
UpdateCTestConfiguration from :/home/lakshmi/work/rpp-test/DartConfiguration.tcl
Parse Config file:/home/lakshmi/work/rpp-test/DartConfiguration.tcl
UpdateCTestConfiguration from :/home/lakshmi/work/rpp-test/DartConfiguration.tcl
Parse Config file:/home/lakshmi/work/rpp-test/DartConfiguration.tcl
Test project /home/lakshmi/work/rpp-test
Constructing a list of tests
Done constructing a list of tests
Updating test list for fixtures
Added 0 tests to meet fixture requirements
Checking test dependency graph...
Checking test dependency graph end
test 1
Start 1: rpp_sanity_test_brightness_host_f32
1: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/HOST" "/home/lakshmi/work/rpp-test/HOST" "--build-generator" "Unix Makefiles" "--test-command" "Tensor_image_host" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HOST/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test"
1: Working Directory: /home/lakshmi/work/rpp-test
1: Test timeout computed to be: 1500
1: Internal cmake changing into directory: /home/lakshmi/work/rpp-test/HOST
1: ======== CMake output ======
1: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc
1: -- test_suite/HOST: rpp found with find_package(rpp QUIET)
1: -- rpp_INCLUDE_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include
1: -- rpp_LIBRARY -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/librpp.so
1: -- rpp_BACKEND_TYPE -- HIP
1: -- rpp_AUDIO_AUGMENTATIONS_SUPPORT -- OFF
1: -- test_suite/HOST is supported to build and run with RPP HIP backend installation
1: -- Python3 found at /usr/bin/python3.11
1: -- test_suite/HOST All Tensor HOST tests set to build with OpenMP
1: -- test_suite/HOST Tensor-Image HOST tests set to build with OpenCV (decoder_type 1 and image dumps enabled)
1: -- Tensor_image_host -- Link Libraries: see target_link_libraries
1: -- test_suite/HOST Tensor-Voxel HOST tests set to build with libniftiio
1: -- test_suite/HOST Tensor-Voxel HOST tests set to build with OpenCV (slice JPG dumps enabled)
1: -- Warning: RPP_AUDIO_SUPPORT not found in RPP installation. It must be installed to build test_suite/HOST Tensor-Audio tests successfully!
1: Configuring done
1: Generating done
1: Build files have been written to: /home/lakshmi/work/rpp-test/HOST
1: ======== End CMake output ======
1: Change Dir: /home/lakshmi/work/rpp-test/HOST
1:
1: Run Clean Command:/usr/bin/gmake -f Makefile clean
1:
1: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 16%] Building CXX object CMakeFiles/Tensor_image_host.dir/Tensor_image_host.cpp.o
1: [ 33%] Linking CXX executable Tensor_image_host
1: [ 33%] Built target Tensor_image_host
1: [ 50%] Building CXX object CMakeFiles/Tensor_voxel_host.dir/Tensor_voxel_host.cpp.o
1: [ 66%] Linking CXX executable Tensor_voxel_host
1: [ 66%] Built target Tensor_voxel_host
1: [ 83%] Building CXX object CMakeFiles/Tensor_misc_host.dir/Tensor_misc_host.cpp.o
1: [100%] Linking CXX executable Tensor_misc_host
1: [100%] Built target Tensor_misc_host
1:
1: Running test command: "/home/lakshmi/work/rpp-test/HOST/Tensor_image_host" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HOST/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test"
1:
1: Inputs for this test case are:
1: src1 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1
1: src2 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1
1: dst = /home/lakshmi/work/rpp-test/HOST/output
1: u8 / f16 / f32 / u8->f16 / u8->f32 / i8 / u8->i8 (0/1/2/3/4/5/6) = 2
1: outputFormatToggle (pkd->pkd = 0 / pkd->pln = 1) = 0
1: case number (0:101) = 0
1: number of times to run = 1
1: test type - (0 = unit tests / 1 = performance tests) = 0
1: layout type - (0 = PKD3 / 1 = PLN3 / 2 = PLN1) = 0
1: qa mode - 0/1 = 0
1: decoder type - (0 = packed .rgb / 1 = OpenCV) = 0
1: batch size = 1
1: Running brightness_f32_Tensor_HOST_PKD3_toPKD3 1 times (each time with a batch size of 1 images) and computing mean statistics...
1:
1: Execution Timings for Iteration 1:
1: CPU Backend Clock Time: 0.046 ms/batch
1: CPU Backend Wall Time: 0.0491142 ms/batch
1:
1: Execution Timings for Iteration 2:
1: CPU Backend Clock Time: 0.062 ms/batch
1: CPU Backend Wall Time: 0.0629425 ms/batch
1:
1: Execution Timings for Iteration 3:
1: CPU Backend Clock Time: 0.061 ms/batch
1: CPU Backend Wall Time: 0.0619888 ms/batch
1:
1/2 Test #1: rpp_sanity_test_brightness_host_f32 ... Passed 6.16 sec
test 2
Start 2: rpp_sanity_test_brightness_hip_f32
2: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/HIP" "/home/lakshmi/work/rpp-test/HIP" "--build-generator" "Unix Makefiles" "--test-command" "Tensor_image_hip" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HIP/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test"
2: Working Directory: /home/lakshmi/work/rpp-test
2: Test timeout computed to be: 1500
2: Internal cmake changing into directory: /home/lakshmi/work/rpp-test/HIP
2: ======== CMake output ======
2: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc
2: -- test_suite/HIP: rpp found with find_package(rpp QUIET)
2: -- rpp_INCLUDE_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include
2: -- rpp_LIBRARY -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/librpp.so
2: -- rpp_BACKEND_TYPE -- HIP
2: -- rpp_AUDIO_AUGMENTATIONS_SUPPORT -- OFF
2: -- test_suite/HIP is supported to build and run with RPP HIP backend installation
2: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc
2: -- HIP found at HIP_INCLUDE_DIRS - /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include and HIP_LIBRARIES - hip::host;hip::device
2: -- Python3 found at /usr/bin/python3.11
2: -- test_suite/HIP All Tensor HIP tests set to build with OpenMP
2: -- test_suite/HIP Tensor-Image HIP tests set to build with OpenCV (decoder_type 1 and image dumps enabled)
2: -- Tensor_image_hip -- Link Libraries: see target_link_libraries
2: -- test_suite/HIP Tensor-Voxel HIP tests set to build with libniftiio
2: -- test_suite/HIP Tensor-Voxel HIP tests set to build with OpenCV (slice JPG dumps enabled)
2: -- Tensor_voxel_hip -- Link Libraries: see target_link_libraries
2: -- Warning: RPP_AUDIO_SUPPORT not found in RPP installation. It must be installed to build test_suite/HIP Tensor-Audio tests successfully!
2: -- Tensor_misc_hip -- Link Libraries: rpp::rpp;hip::host;hip::device;stdc++fs;OpenMP::OpenMP_CXX
2: Configuring done
2: Generating done
2: Build files have been written to: /home/lakshmi/work/rpp-test/HIP
2: ======== End CMake output ======
2: Change Dir: /home/lakshmi/work/rpp-test/HIP
2:
2: Run Clean Command:/usr/bin/gmake -f Makefile clean
2:
2: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 16%] Building CXX object CMakeFiles/Tensor_image_hip.dir/Tensor_image_hip.cpp.o
2: [ 33%] Linking CXX executable Tensor_image_hip
2: [ 33%] Built target Tensor_image_hip
2: [ 50%] Building CXX object CMakeFiles/Tensor_voxel_hip.dir/Tensor_voxel_hip.cpp.o
2: [ 66%] Linking CXX executable Tensor_voxel_hip
2: [ 66%] Built target Tensor_voxel_hip
2: [ 83%] Building CXX object CMakeFiles/Tensor_misc_hip.dir/Tensor_misc_hip.cpp.o
2: [100%] Linking CXX executable Tensor_misc_hip
2: [100%] Built target Tensor_misc_hip
2:
2: Running test command: "/home/lakshmi/work/rpp-test/HIP/Tensor_image_hip" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HIP/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test"
2:
2: Inputs for this test case are:
2: src1 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1
2: src2 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1
2: dst = /home/lakshmi/work/rpp-test/HIP/output
2: u8 / f16 / f32 / u8->f16 / u8->f32 / i8 / u8->i8 (0/1/2/3/4/5/6) = 2
2: outputFormatToggle (pkd->pkd = 0 / pkd->pln = 1) = 0
2: case number (0:101) = 0
2: number of times to run = 1
2: test type - (0 = unit tests / 1 = performance tests) = 0
2: layout type - (0 = PKD3/ 1 = PLN3/ 2 = PLN1) = 0
2: qa mode - 0/1 = 0
2: decoder type - (0 = packed .rgb / 1 = OpenCV) = 0
2: batch size = 1
2: Running brightness_f32_Tensor_HIP_PKD3_toPKD3 1 times (each time with a batch size of 1 images) and computing mean statistics...
2:
2: Execution Timings for Iteration 1:
2: GPU Backend Wall Time: 3.59392 ms/batch
2:
2: Execution Timings for Iteration 2:
2: GPU Backend Wall Time: 0.0441074 ms/batch
2:
2: Execution Timings for Iteration 3:
2: GPU Backend Wall Time: 0.0441074 ms/batch
2/2 Test ROCm#2: rpp_sanity_test_brightness_hip_f32 .... Passed 13.04 sec
100% tests passed, 0 tests failed out of 2
Total Test time (real) = 19.20 sec
```
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
goldcoderZ
pushed a commit
that referenced
this pull request
Jul 22, 2026
…hipcxx (#1) * Remove duplicates from inc/gpu/__algorithm * Remove inc/gpu/__atomic * Remove __memory duplicates * Update __memory usages * Remove __iterator duplicates * Remove __numeric duplicates * Remove __type_traits duplicates * Remove __functional duplicates * Update __functional usage in __thread/worknode.h * Remove cstring duplicate * Remove cstdlib duplicate * Remove __tuple_dir duplicate * Remove __utility duplicates * Update __utility usage * Remove __split_buffer header * Remove vector header * Remove memcmp.h from __clib * Remove __string dir * Remove device unique_ptr tests * Remove host unique_ptr tests * Remove exposed unique_ptr_test.cxx * Remove vector_test * Update compile_test.cxx * Remove atomic_test.cxx * Remove atomic lit tests --------- Co-authored-by: Marko Savic <marsavic@amd.com>
goldcoderZ
pushed a commit
that referenced
this pull request
Jul 22, 2026
…hipcxx (#1) * Remove duplicates from inc/gpu/__algorithm * Remove inc/gpu/__atomic * Remove __memory duplicates * Update __memory usages * Remove __iterator duplicates * Remove __numeric duplicates * Remove __type_traits duplicates * Remove __functional duplicates * Update __functional usage in __thread/worknode.h * Remove cstring duplicate * Remove cstdlib duplicate * Remove __tuple_dir duplicate * Remove __utility duplicates * Update __utility usage * Remove __split_buffer header * Remove vector header * Remove memcmp.h from __clib * Remove __string dir * Remove device unique_ptr tests * Remove host unique_ptr tests * Remove exposed unique_ptr_test.cxx * Remove vector_test * Update compile_test.cxx * Remove atomic_test.cxx * Remove atomic lit tests --------- Co-authored-by: Marko Savic <marsavic@amd.com> [ROCm/hipthreads commit: 6e7b1eb]
goldcoderZ
pushed a commit
that referenced
this pull request
Jul 22, 2026
## What
Make origami's installed `CTestTestfile.cmake` use paths relative to the
test
directory instead of `${CMAKE_CURRENT_LIST_DIR}`.
## Why
`${CMAKE_CURRENT_LIST_DIR}` is **empty** when `ctest` parses an
*installed*
`CTestTestfile.cmake` — its restricted parser does not populate that
variable
(unlike a normal configure or `cmake -P` run). The relocatable file
emitted by
ROCm#9047 used it for the test command, `WORKING_DIRECTORY`, and the
`LD_LIBRARY_PATH`/`PYTHONPATH` environment, so `WORKING_DIRECTORY`
expanded to
`""` and every test failed to launch:
```
Failed to change working directory to : No such file or directory
1/2 Test #1: origami-tests ....................***Not Run
2/2 Test ROCm#2: origami_python_tests .............***Not Run
```
ctest already runs each test with its working directory defaulting to
the
directory containing the `CTestTestfile.cmake`, including under
`ctest --test-dir <dir>` invoked from an unrelated cwd. So relative
paths
resolve correctly and relocatably without the (empty) variable. This
also aligns
origami with the sibling hipBLASLt installed test file, which uses the
same
relative-path convention ("Tests are defined with relative paths to work
in the
installed location").
## Change
`shared/origami/tests/CMakeLists.txt` — in the installed
`_ctest_content`:
- `"${CMAKE_CURRENT_LIST_DIR}/../origami-tests"` → `"../origami-tests"`
- pytest target `"${CMAKE_CURRENT_LIST_DIR}/tests"` → `"tests"`
- dropped `WORKING_DIRECTORY "${CMAKE_CURRENT_LIST_DIR}"` (default is
already the file's dir)
- env paths `${CMAKE_CURRENT_LIST_DIR}/../../lib[...]` →
`../../lib[...]`
## Verification
Reproduced and validated in
`ghcr.io/rocm/therock_build_manylinux_x86_64`
(ctest 3.27.9, python 3.12) against a real origami build. Both installed
`CTestTestfile.cmake` variants were generated by actual `cmake`
CONFIGURE from
the committed source and run via `ctest --test-dir` from an unrelated
cwd:
| | C++ `origami-tests` | `origami_python_tests` |
|---|---|---|
| Before (`${CMAKE_CURRENT_LIST_DIR}`) | Not Run | Not Run |
| After (relative) | Passed | Runs |
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
goldcoderZ
pushed a commit
that referenced
this pull request
Jul 22, 2026
… uchar type (ROCm#9237) ## Motivation This PR addresses compile warnings with Ofast flag. Instead changed to use -fast-math ## Technical Details Modified compile flags to use standard O3 Removed the use of uchar data-type from RPP test-suite which is only defined in OpenCV ## Test Plan CTests should pass ## Test Result Start 1: rpp_sanity_test_brightness_host_f32 1/10 Test #1: rpp_sanity_test_brightness_host_f32 ......... Passed 3.64 sec Start 2: rpp_sanity_test_brightness_hip_f32 2/10 Test ROCm#2: rpp_sanity_test_brightness_hip_f32 .......... Passed 7.75 sec Start 3: rpp_qa_tests_tensor_image_host_all 3/10 Test ROCm#3: rpp_qa_tests_tensor_image_host_all .......... Passed 42.43 sec Start 4: rpp_qa_tests_tensor_voxel_host_all 4/10 Test ROCm#4: rpp_qa_tests_tensor_voxel_host_all .......... Passed 2.26 sec Start 5: rpp_qa_tests_tensor_misc_host_all 5/10 Test ROCm#5: rpp_qa_tests_tensor_misc_host_all ........... Passed 7.74 sec Start 6: rpp_qa_tests_tensor_misc_host_test_type_1 6/10 Test ROCm#6: rpp_qa_tests_tensor_misc_host_test_type_1 ... Passed 402.60 sec Start 7: rpp_qa_tests_tensor_image_hip_all 7/10 Test ROCm#7: rpp_qa_tests_tensor_image_hip_all ........... Passed 84.44 sec Start 8: rpp_qa_tests_tensor_voxel_hip_all 8/10 Test ROCm#8: rpp_qa_tests_tensor_voxel_hip_all ........... Passed 4.62 sec Start 9: rpp_qa_tests_tensor_misc_hip_all 9/10 Test ROCm#9: rpp_qa_tests_tensor_misc_hip_all ............ Passed 16.50 sec Start 10: rpp_qa_tests_tensor_misc_hip_test_type_1 10/10 Test ROCm#10: rpp_qa_tests_tensor_misc_hip_test_type_1 .... Passed 348.19 sec 100% tests passed, 0 tests failed out of 10 ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
goldcoderZ
pushed a commit
that referenced
this pull request
Sep 27, 2026
…nitialization (ROCm#9977) ## Motivation This PR addresses issue [738](ROCm/rpp#738) in ROCm/rpp. In `rpp/src/modules/handle_hip.cpp`, `PreInitializeBuffer()` allocates a new `InitHandle` object and then calls `PreInitializeBufferCPU()`, which allocates that object again. This causes a (albeit small) memory leak of the first `InitHandle` object. ## Technical Details Remove the first allocation (`this->initHandle = new InitHandle();`) in `PreInitializeBuffer()`, since it will always call `PreInitializeBufferCPU()` immediately after. This way there is only one allocation in total. ## Test Plan CTests should pass. ## Test Result ``` Start 1: rpp_sanity_test_brightness_host_f32 1/2 Test #1: rpp_sanity_test_brightness_host_f32 ... Passed 3.13 sec Start 2: rpp_sanity_test_brightness_hip_f32 2/2 Test ROCm#2: rpp_sanity_test_brightness_hip_f32 .... Passed 6.36 sec ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>
goldcoderZ
pushed a commit
that referenced
this pull request
Sep 27, 2026
## Motivation As per https://github.com/ROCm/TheRock/blob/main/docs/development/test_filtering.md, we want to add test filter options to RPP. ## JIRA ID ROCM-26141 ## Technical Details This PR adds these options to rpp ctests, with the existing test base. More tests can be added as required. | Label | Tests | Target Time | |-----------------|--------------------------------------------------------------------------------------------|-------------| | quick | rpp_sanity_test_brightness_host_f32, rpp_sanity_test_brightness_hip_f32 | < 5 min | | standard | quick tests + rpp_qa_tests_tensor_misc_host_all, rpp_qa_tests_tensor_misc_hip_all | < 30 min | | comprehensive | All rpp_qa_tests_tensor_*_all (image, voxel, audio, misc for both HOST and HIP) | < 2 hours | | full | Everything including qa_mode_0 and test_type_1 (performance) tests | hours+ | ## Test Plan Test ctests with these labels: ctest -L quick # Sanity checks only ctest -L standard # PR validation ctest -L comprehensive # Nightly tests ctest -L full # Full test suite ## Test Result Example of quick test results: ``` lakshmi@kapu:~/work/rpp-test$ ctest -VV -L quick UpdateCTestConfiguration from :/home/lakshmi/work/rpp-test/DartConfiguration.tcl Parse Config file:/home/lakshmi/work/rpp-test/DartConfiguration.tcl UpdateCTestConfiguration from :/home/lakshmi/work/rpp-test/DartConfiguration.tcl Parse Config file:/home/lakshmi/work/rpp-test/DartConfiguration.tcl Test project /home/lakshmi/work/rpp-test Constructing a list of tests Done constructing a list of tests Updating test list for fixtures Added 0 tests to meet fixture requirements Checking test dependency graph... Checking test dependency graph end test 1 Start 1: rpp_sanity_test_brightness_host_f32 1: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/HOST" "/home/lakshmi/work/rpp-test/HOST" "--build-generator" "Unix Makefiles" "--test-command" "Tensor_image_host" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HOST/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test" 1: Working Directory: /home/lakshmi/work/rpp-test 1: Test timeout computed to be: 1500 1: Internal cmake changing into directory: /home/lakshmi/work/rpp-test/HOST 1: ======== CMake output ====== 1: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 1: -- test_suite/HOST: rpp found with find_package(rpp QUIET) 1: -- rpp_INCLUDE_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include 1: -- rpp_LIBRARY -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/librpp.so 1: -- rpp_BACKEND_TYPE -- HIP 1: -- rpp_AUDIO_AUGMENTATIONS_SUPPORT -- OFF 1: -- test_suite/HOST is supported to build and run with RPP HIP backend installation 1: -- Python3 found at /usr/bin/python3.11 1: -- test_suite/HOST All Tensor HOST tests set to build with OpenMP 1: -- test_suite/HOST Tensor-Image HOST tests set to build with OpenCV (decoder_type 1 and image dumps enabled) 1: -- Tensor_image_host -- Link Libraries: see target_link_libraries 1: -- test_suite/HOST Tensor-Voxel HOST tests set to build with libniftiio 1: -- test_suite/HOST Tensor-Voxel HOST tests set to build with OpenCV (slice JPG dumps enabled) 1: -- Warning: RPP_AUDIO_SUPPORT not found in RPP installation. It must be installed to build test_suite/HOST Tensor-Audio tests successfully! 1: Configuring done 1: Generating done 1: Build files have been written to: /home/lakshmi/work/rpp-test/HOST 1: ======== End CMake output ====== 1: Change Dir: /home/lakshmi/work/rpp-test/HOST 1: 1: Run Clean Command:/usr/bin/gmake -f Makefile clean 1: 1: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 16%] Building CXX object CMakeFiles/Tensor_image_host.dir/Tensor_image_host.cpp.o 1: [ 33%] Linking CXX executable Tensor_image_host 1: [ 33%] Built target Tensor_image_host 1: [ 50%] Building CXX object CMakeFiles/Tensor_voxel_host.dir/Tensor_voxel_host.cpp.o 1: [ 66%] Linking CXX executable Tensor_voxel_host 1: [ 66%] Built target Tensor_voxel_host 1: [ 83%] Building CXX object CMakeFiles/Tensor_misc_host.dir/Tensor_misc_host.cpp.o 1: [100%] Linking CXX executable Tensor_misc_host 1: [100%] Built target Tensor_misc_host 1: 1: Running test command: "/home/lakshmi/work/rpp-test/HOST/Tensor_image_host" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HOST/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test" 1: 1: Inputs for this test case are: 1: src1 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1 1: src2 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1 1: dst = /home/lakshmi/work/rpp-test/HOST/output 1: u8 / f16 / f32 / u8->f16 / u8->f32 / i8 / u8->i8 (0/1/2/3/4/5/6) = 2 1: outputFormatToggle (pkd->pkd = 0 / pkd->pln = 1) = 0 1: case number (0:101) = 0 1: number of times to run = 1 1: test type - (0 = unit tests / 1 = performance tests) = 0 1: layout type - (0 = PKD3 / 1 = PLN3 / 2 = PLN1) = 0 1: qa mode - 0/1 = 0 1: decoder type - (0 = packed .rgb / 1 = OpenCV) = 0 1: batch size = 1 1: Running brightness_f32_Tensor_HOST_PKD3_toPKD3 1 times (each time with a batch size of 1 images) and computing mean statistics... 1: 1: Execution Timings for Iteration 1: 1: CPU Backend Clock Time: 0.046 ms/batch 1: CPU Backend Wall Time: 0.0500679 ms/batch 1: 1: Execution Timings for Iteration 2: 1: CPU Backend Clock Time: 0.06 ms/batch 1: CPU Backend Wall Time: 0.0619888 ms/batch 1: 1: Execution Timings for Iteration 3: 1: CPU Backend Clock Time: 0.062 ms/batch 1: CPU Backend Wall Time: 0.0629425 ms/batch 1: 1/2 Test #1: rpp_sanity_test_brightness_host_f32 ... Passed 6.14 sec test 2 Start 2: rpp_sanity_test_brightness_hip_f32 2: Test command: /usr/local/bin/ctest "--build-and-test" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/HIP" "/home/lakshmi/work/rpp-test/HIP" "--build-generator" "Unix Makefiles" "--test-command" "Tensor_image_hip" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HIP/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test" 2: Working Directory: /home/lakshmi/work/rpp-test 2: Test timeout computed to be: 1500 2: Internal cmake changing into directory: /home/lakshmi/work/rpp-test/HIP 2: ======== CMake output ====== 2: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 2: -- test_suite/HIP: rpp found with find_package(rpp QUIET) 2: -- rpp_INCLUDE_DIR -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include 2: -- rpp_LIBRARY -- /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/lib/librpp.so 2: -- rpp_BACKEND_TYPE -- HIP 2: -- rpp_AUDIO_AUGMENTATIONS_SUPPORT -- OFF 2: -- test_suite/HIP is supported to build and run with RPP HIP backend installation 2: HIP: Using hipcc from relative path: /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/bin/hipcc 2: -- HIP found at HIP_INCLUDE_DIRS - /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/include and HIP_LIBRARIES - hip::host;hip::device 2: -- Python3 found at /usr/bin/python3.11 2: -- test_suite/HIP All Tensor HIP tests set to build with OpenMP 2: -- test_suite/HIP Tensor-Image HIP tests set to build with OpenCV (decoder_type 1 and image dumps enabled) 2: -- Tensor_image_hip -- Link Libraries: see target_link_libraries 2: -- test_suite/HIP Tensor-Voxel HIP tests set to build with libniftiio 2: -- test_suite/HIP Tensor-Voxel HIP tests set to build with OpenCV (slice JPG dumps enabled) 2: -- Tensor_voxel_hip -- Link Libraries: see target_link_libraries 2: -- Warning: RPP_AUDIO_SUPPORT not found in RPP installation. It must be installed to build test_suite/HIP Tensor-Audio tests successfully! 2: -- Tensor_misc_hip -- Link Libraries: rpp::rpp;hip::host;hip::device;stdc++fs;OpenMP::OpenMP_CXX 2: Configuring done 2: Generating done 2: Build files have been written to: /home/lakshmi/work/rpp-test/HIP 2: ======== End CMake output ====== 2: Change Dir: /home/lakshmi/work/rpp-test/HIP 2: 2: Run Clean Command:/usr/bin/gmake -f Makefile clean 2: 2: Run Build Command(s):/usr/bin/gmake -f Makefile && [ 16%] Building CXX object CMakeFiles/Tensor_image_hip.dir/Tensor_image_hip.cpp.o 2: [ 33%] Linking CXX executable Tensor_image_hip 2: [ 33%] Built target Tensor_image_hip 2: [ 50%] Building CXX object CMakeFiles/Tensor_voxel_hip.dir/Tensor_voxel_hip.cpp.o 2: [ 66%] Linking CXX executable Tensor_voxel_hip 2: [ 66%] Built target Tensor_voxel_hip 2: [ 83%] Building CXX object CMakeFiles/Tensor_misc_hip.dir/Tensor_misc_hip.cpp.o 2: [100%] Linking CXX executable Tensor_misc_hip 2: [100%] Built target Tensor_misc_hip 2: 2: Running test command: "/home/lakshmi/work/rpp-test/HIP/Tensor_image_hip" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1" "/home/lakshmi/work/rpp-test/HIP/output" "2" "0" "0" "0" "1" "0" "0" "1" "0" "0" "1" "0" "0" "0" "0" "/home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test" 2: 2: Inputs for this test case are: 2: src1 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1 2: src2 = /home/lakshmi/rock_0408/lib/python3.10/site-packages/_rocm_sdk_devel/share/rpp/test/TEST_IMAGES/three_images_mixed_src1 2: dst = /home/lakshmi/work/rpp-test/HIP/output 2: u8 / f16 / f32 / u8->f16 / u8->f32 / i8 / u8->i8 (0/1/2/3/4/5/6) = 2 2: outputFormatToggle (pkd->pkd = 0 / pkd->pln = 1) = 0 2: case number (0:101) = 0 2: number of times to run = 1 2: test type - (0 = unit tests / 1 = performance tests) = 0 2: layout type - (0 = PKD3/ 1 = PLN3/ 2 = PLN1) = 0 2: qa mode - 0/1 = 0 2: decoder type - (0 = packed .rgb / 1 = OpenCV) = 0 2: batch size = 1 2: Running brightness_f32_Tensor_HIP_PKD3_toPKD3 1 times (each time with a batch size of 1 images) and computing mean statistics... 2: 2: Execution Timings for Iteration 1: 2: GPU Backend Wall Time: 3.73983 ms/batch 2: 2: Execution Timings for Iteration 2: 2: GPU Backend Wall Time: 0.0448227 ms/batch 2: 2: Execution Timings for Iteration 3: 2: GPU Backend Wall Time: 0.0429153 ms/batch 2/2 Test ROCm#2: rpp_sanity_test_brightness_hip_f32 .... Passed 13.19 sec 100% tests passed, 0 tests failed out of 2 Label Time Summary: all = 19.33 sec*proc (2 tests) comprehensive = 19.33 sec*proc (2 tests) full = 19.33 sec*proc (2 tests) nightly = 19.33 sec*proc (2 tests) pr = 19.33 sec*proc (2 tests) pre-commit = 19.33 sec*proc (2 tests) quick = 19.33 sec*proc (2 tests) smoke = 19.33 sec*proc (2 tests) standard = 19.33 sec*proc (2 tests) weekly = 19.33 sec*proc (2 tests) Total Test time (real) = 19.33 sec ``` ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Kiriti Gowda <kiritigowda@gmail.com>
goldcoderZ
pushed a commit
that referenced
this pull request
Sep 27, 2026
ISSUE ID: ROCm#8997 ## Motivation The TileEngine → Dispatcher bridge had no path for the weight-preshuffled GEMM op **gemm_preshuffle**, which pre-permutes the B (weight) tensor into the pipeline's packed layout for higher throughput on weight-heavy GEMMs. This is a real Old-TE capability with no dispatcher equivalent, so this PR adds the bridge so the dispatcher can generate and launch preshuffle GEMM at parity with the legacy Tile Engine version. The capability set matches the Old-TE op exactly: `fp16`, `bf16`, `fp8`, `bf8`, `rcr` layout only (A row-major, B col-major, C row-major), with the B tensor preshuffled into the packed layout. It follows the same bridge scheme as ROCm#8997 (regular GEMM), ROCm#9000 (grouped), and ROCm#9028 (stream-K). ## Test Plan - Run the CPU-only unit tests (no GPU required): `python3 -m pytest dispatcher/tests/test_preshuffle_bridge.py -v` - Codegen + build the `.so` for all 4 dtypes and verify output vs an fp32 NumPy reference at M=N=K=512 on gfx942 / MI300X. - Confirm the preshuffle B-shuffle matches `ck_tile::shuffle_b` byte-for-byte. ## Test Result - CPU-only unit tests pass (8 passed). - Codegen + `.so` build succeed for all 4 dtypes (0 codegen failures). Output verified vs fp32 reference at M=N=K=512: | dtype | max_rel | status | |-------|---------|--------| | fp16 | 5.2e-4 | PASS | | bf16 | 5.5e-3 | PASS | | fp8 | 3.4e-2 | PASS | | bf8 | 6.5e-2 | PASS | - The preshuffle B-shuffle (the #1 correctness risk) is verified byte-for-byte via `ck_tile::shuffle_b`; the permuteN shuffle gives ~1.25 rel error and is deliberately not used, since the bridge does not emit the permuteN pipeline. > **Note (arch support):** the correctness above is on gfx942 (MI300X). fp8/bf8 preshuffle is unsupported on **gfx950** for *both* this bridge and Old-TE — the underlying kernel fails to launch there (`status -2`) — while fp16/bf16 are correct on both arches (see the cross-arch note in the comments). --- **Related PRs / references (TileEngine → Dispatcher GEMM bridge series):** ROCm#8997 (regular GEMM fp16/bf16 all-layout), ROCm#9000 (grouped GEMM), ROCm#9028 (stream-K), ROCm#8887 (fp8/bf8/int8). This PR is a sibling in the same bridge effort tracked across those PRs. --------- Co-authored-by: Muhammed Emin Ozturk <3836908+ozturkosu@users.noreply.github.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> Co-authored-by: Thrupti Raj Lakshmana Gowda <thruptiraj.lakshmanagowda@amd.com>
goldcoderZ
pushed a commit
that referenced
this pull request
Sep 27, 2026
…Cm#11978) ## Summary Completes the library carve for convolution, following the attention split (ROCm#8977). Every convolution kernel builder, dispatch policy, test, parity emitter, benchmark, and builder harness now lives under `rocke/library/`. `rocke/platform/` keeps none of them — not even re-export shims — except where a real architectural constraint (the manifest-runner's lazy-resolution registry) requires a thin pointer. Also ports the two-stage deterministic wgrad feature (ROCm#10571) into the library, since that PR landed on `develop` after this branch was cut, and fixes a handful of bugs found in review (see "Bug fixes" below). This is a full move, not a copy: `git diff --diff-filter=D` against the merge-base shows the platform originals gone, not left behind. An earlier version of this PR left the platform copies in place as unreferenced duplicates; that duplication has been removed in this revision. JIRA-ID: AICK-2047 --- ## What changed and why ### Kernel builders (`library/kernels/common/` and per-arch subpackages) | File | Description | |---|---| | `conv_implicit_gemm.py` | Forward NHWC implicit-GEMM | | `conv_implicit_gemm_wgrad.py` | Backward-weight implicit-GEMM; adds `two_stage` / `force_deterministic` fields + workspace-store epilogue (ported from ROCm#10571) | | `conv_implicit_gemm_wgrad_two_stage.py` | Two-stage pipeline launcher (ported from ROCm#10571) | | `conv_wgrad_workspace_reduce.py` | Stage 2 workspace-reduce kernel builder (ported from ROCm#10571) | | `conv_implicit_gemm_dgrad.py` | Backward-data implicit-GEMM | | `conv_direct_grouped.py` | Direct grouped conv (4c / 8c / 16c / 32c / depthwise variants) | | `deep_fused_conv_pool.py` + gfx950 / gfx1151 / gfx1201 arch variants | Fused conv+pool | | `img2col.py` | Im2col transform | All `build_*` functions take `arch` as a keyword-only parameter, per the builder contract established in ROCm#11237. ### Manifest runner `platform/python/rocke/run_manifest.py` no longer imports the conv manifest runner eagerly at module scope. It resolves `conv_fp16` / `conv_bf16` / `conv_fp32` lazily through the existing `_LIBRARY_RUNNER_MODULES` registry (the same pattern already used for `deep_fused_conv_pool_fp16` / `_i8i4`), pointing at `kernels.common.manifest_runner.conv`. This is the one place `platform` still names a `library` module path, and it does so only inside a function body, resolved at call time — `platform` stays standalone-installable with `library` absent. ### Tests | File | Description | |---|---| | `library/tests/test_conv_fwd_correctness.py` | GPU correctness for fwd | | `library/tests/test_conv_wgrad_correctness.py` | GPU correctness for wgrad; includes `TestConvWgradTwoStage` (29 methods, ported from ROCm#10571) | | `library/tests/test_conv_dgrad_correctness.py` | GPU correctness for dgrad | | `library/tests/test_direct_conv_correctness.py` | GPU correctness for direct-grouped conv (moved from `platform/tests/instances/`) | | `library/tests/test_conv_build_signatures.py` / `test_builder_signature_contract.py` | Build-signature contracts (no GPU) | | `library/tests/parity/conv_implicit_gemm_{,wgrad,dgrad}_emit.{py,c}` | Python/C byte-identity emitters; wgrad extended with the two-stage and split-K configs ported from ROCm#10571 | | `library/tests/parity/conv_wgrad_workspace_reduce_emit.{py,c}` | Parity configs for the Stage 2 reduce kernel | | `library/tests/parity/conv_direct_grouped_emit.{py,c}` | Extended to cover all config variants (16c/4c/8c/32c/depthwise) that platform's copy had and library's did not | | `library/tests/parity/img2col_emit.{py,c}` | Moved from `platform/tests/instances/parity/` — it already imported `kernels.common.*` at module scope and had no library-side counterpart | | `library/tests/dispatch/test_grouped_conv_wgrad_dispatch.py` | Wgrad dispatch selection | `library/tests/test_conv_dgrad_correctness.py` resolves `platform`/`library` roots via `rocke.assets.library_root()` / `platform_root()` instead of hand-rolled `parents[N]` math (the old math pointed one level too high and silently ran its subprocess benchmark with a broken `PYTHONPATH`). The child process's own `PYTHONPATH` is now built by prefixing the current process's `sys.path`, so the same test works unchanged whether it's running from a source checkout or an installed CTest artifact. ### Benchmarks / builders `benchmark_implicit_gemm_conv.py`, `benchmark_direct_conv.py`, `conv_reference.py`, `benchmark_conv_compare.py`, and the gfx950 / gfx1151 / gfx1201 deep-conv-fusion builder harnesses all live under `library/benchmarks/common/` and `library/builders/<arch>/` now. `benchmark_implicit_gemm_conv.py` includes the `--two-stage` sweep support (CLI flag, Stage 1 + Stage 2 GPU run, results reporting) that had been left out of an earlier draft of this move. `library/builders/common/conv_reference.py` now slices the weight tensor to `C // groups` channels before the torch reference conv, matching platform's original grouped-conv fix — without it, reference values for `groups > 1` were wrong. ### Documentation Fixed every dangling reference to the old `platform/python/rocke/instances/common/` and `platform/tests/instances/` conv paths across `dsl_docs/`, builder `README.md`s, and in-code docstrings (Run: instructions, "Key Files" tables, prose) — both broken markdown links and plain-text mentions. ### CI - `rocke_library_conv_pytest` (existing CTest target): CPU-only dispatch/signature tests, no GPU required. - `rocke_library_conv_gpu_pytest` (new): registers `test_conv_{fwd,wgrad,dgrad}_correctness.py` and `test_direct_conv_correctness.py`. These were collected by the platform-wide `rocke_pytest` target while they lived under `platform/tests/instances/`; moving them to `library/tests/` put them inside that target's `--ignore=tests/library` scope, silently dropping them from CI. They self-skip when no supported GPU/torch is detected, so a host-only runner still passes; the `gpu` label lets a runner select or exclude them explicitly. GPU-nightly runner registration (which machines actually run the `gpu`-labelled set) is out of scope here — no such runner config exists anywhere in this repo for me to safely edit; that needs a human with access to that infra. ### Bug fixes found during review and re-verification - `platform/python/rocke/core/lower_cktile.py` deferred its `kernels.common.conv_*` import until after the `UniversalGemmSpec` branch — it was unconditionally importing a `library`-layer module before checking which spec type it had, breaking plain-GEMM lowering on a `platform`-only install. - `platform/python/rocke/examples/common/hip_lowering_parity.py` had a module-level `kernels.common.*` import that broke standalone-platform installs for all ~40 non-conv families this audit tool also covers; deferred into the function body that actually needs it. - Both `test_direct_conv_correctness.py` and `test_conv_fwd_correctness.py` had a pre-existing HIP-context race (predates this migration, confirmed via `git blame` against ROCm#11168): calling rocke's HIP runtime before torch touches CUDA claims the process HIP context and permanently breaks `torch.cuda()` for the rest of the process. Fixed by having torch claim the context first when present. - Deleted a stale `platform/python/rocke/examples/gfx1151/deep_conv_fusion/` copy whose `compare_prebuilt.py` imported a `run_manifest` symbol that no longer exists (removed when the manifest runner was made lazy); the library copy was already the correct, current one. --- ## Verification - **Byte-identity gate** (the repo's #1 invariant): GREEN at both `llvm20` and `llvm22`, all families, after every kernel-file move. - **Platform pytest** (standalone, `library` absent from `sys.path`): 853 passed, 64 skipped, 1548 subtests passed. - **Library CPU tests** (layering, build-signature contract, dispatch): 326 passed, 216 subtests passed. - **GPU correctness, gfx942 (MI300X), run locally in this environment**: `test_conv_fwd_correctness.py`, `test_conv_wgrad_correctness.py`, `test_conv_dgrad_correctness.py`, `test_direct_conv_correctness.py` together — 73 passed, 68 skipped (arch-gated: gfx950/gfx1250-only variants not exercisable on this hardware), 366 subtests passed, ~5.5 minutes. - **Standalone-install invariant**: verified with a script that strips `library` from `sys.path` and imports every `rocke.*` submodule via `pkgutil.walk_packages` — zero failures. - **No architecture violations**: `grep -rn "rocke\.instances" library/` returns only pre-existing, unrelated attention/matmul_nbits references; zero `kernels`/`builders`/`dispatch` imports at module scope anywhere under `platform/python/`. ## Known, pre-existing gaps (not introduced or hidden by this PR) - **No dgrad dispatch policy layer.** Unlike fwd/wgrad, there is no `dispatch/conv_dgrad.py` request-routing module, in either tree, with any git history — dgrad is built directly today. Writing one needs real dispatch-policy design input; out of scope here rather than invented blind on a shipping path. - **gfx950 hardware was not available in this environment.** gfx950-specific conv paths (deep-fused-conv-pool, WMMA variants) are byte-identity-checked but not GPU-run here; they need a human with gfx950 access to sign off numerically. --------- Co-authored-by: Claude <noreply@anthropic.com>
goldcoderZ
pushed a commit
that referenced
this pull request
Sep 27, 2026
… tests (ROCm#12492) ISSUE ID : ROCm#12491 ## Motivation `Tensile/Tests/common` (real codegen, build, and execution against a CPU reference) only runs in Math CI's `preliminary` job today. It doesn't run anywhere in TheRock Multi-Arch CI or nightly. TESTING.md already lists this as a known gap (Improvement Roadmap item 5). This PR adds the category TheRock needs to run it. The TheRock side is a companion PR, linked below. ## Technical Details - Adds a `hw-common` category to `projects/hipblaslt/tensilelite/test_categories.yaml`. It runs `Tensile/Tests/common` with `-m common` and the prebuilt `tensilelite-client` from the artifact. - Arch selection is opt-out, the same way `preliminary` does it. `configMarks()` tags every config `common` and turns each config's `skip-gfxNNNN` marks into skips for whatever arch `rocm_agent_enumerator` reports. So this category deliberately doesn't use `{AMDGPU_TARGETS}` (that would match almost nothing on real hardware), and it doesn't pass `--gpu-targets`. Hardware detection also sets `PyTestBuildArchNames`, which is what lets a gfx120X runner build for whichever of gfx1200/gfx1201 it actually has. - `parallel_workers: 4`, which is what Math CI runs. I didn't copy the 16 from the FFM categories. `assign_gpu_to_worker()` maps workers with `worker_num % num_gpus` and hands out per-worker locks, so on TheRock's 1-GPU runners every worker shares GPU 0. - Per-test timeout is 5400s, as a hang guard only. Math CI runs this suite with no per-test timeout, and the slowest config I found in recent `preliminary` logs (`streamk/sk_I8gemm_quick.yaml`) took 3173s. - Updates `tensilelite/TESTING.md`: the CI-gates section, the TheRock row in the lanes table (it still pointed at the retired `test_tensilelite.py`), two Known Risks rows, and roadmap item 5 (now the deferred sharding follow-up). - Worker-to-GPU assignment on TheRock's 1-GPU runners is fixed by ROCm#11898 (merged separately), which this PR now picks up from develop. The first TheRock run of `hw-common` hit exactly that bug: all 76 failures were on xdist workers gw1-gw3 with `no ROCm-capable device is detected`. With a worker-assignment fix in place, a later run passed 135 of 136. The one failure was `gemm/xfp32.yaml`, which ROCm#12569 fixes. Nothing selects `hw-common` until the TheRock PR lands, so this is inert on its own. ## Test Plan - Loaded the YAML and ran TheRock's `pytest_runner.build_marker_expression()` against `hw-common` for each target family, to confirm the `-m` expression it builds. - Sized the timeouts from four recent green `preliminary` builds (PR-12472 ROCm#2, PR-12469 #1, PR-12462 #1, PR-12461 #1). - Real collection and execution need a HIP runtime, so that happens on the TheRock PR's GPU run. ## Test Result - Marker expressions come out as `(common) and not skip-<arch>` for gfx90a, gfx94X, gfx950, and gfx120X. - Measured `-m common` wall clock in `preliminary` at 4 workers: gfx950 1:43 to 2:03 (182 tests), gfx942 49 to 64 min (136), gfx90a 16 to 24 min (43), gfx12 13 to 25 min (38). - Ran `hw-common` on gfx942 through TheRock Multi-Arch CI ([run 35915539444](https://github.com/ROCm/rocm-libraries/actions/runs/35915539444)): 136 tests selected, matching Math CI's `preliminary`. Every test on gw0 passed (46/46). All 76 failures were on gw1-gw3 with `no ROCm-capable device is detected`, which is what the conftest fix addresses. The suite ran in 35 minutes. - With ROCm#12569 (the msgpack LibraryClient fix) below it in the stack and ROCm#12526's TheRock pins on top, `tensilelite-common` passed all 136 tests on gfx942 in 46 minutes on the stack's testing PR, ROCm#12573 ([job](https://github.com/ROCm/rocm-libraries/actions/runs/36062518748/job/107875479911)). That includes `gemm/xfp32.yaml`, the one test that failed before ROCm#12569. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. It's a new category plus docs, and nothing selects it until the TheRock PR merges. Related: ROCm/TheRock#8459 (companion), ROCm/TheRock#5183, ROCm/TheRock#3288 Co-authored-by: Cursor <cursoragent@cursor.com>
goldcoderZ
pushed a commit
that referenced
this pull request
Sep 27, 2026
# Community Sample: Flash Attention v2 Forward (`simple_flash_attn_v2.cpp`) ## Motivation This PR adds a new community sample that demonstrates a **Flash Attention v2 style forward kernel** using the rocWMMA API. The sample implements the attention forward operation in a single fused GPU kernel: ```cpp O = softmax(Q * K^T * scale) * V ``` where `scale = 1 / sqrt(D)` and the head dimension `D` is fixed to 64 in this sample. Flash Attention is the standard memory-efficient attention algorithm used in modern LLM inference and training stacks. Instead of materializing the full `[N x N]` attention score matrix in global memory, the kernel processes K/V tiles incrementally and keeps the running softmax statistics in registers / LDS: ```cpp m_i = running row max l_i = running exp-sum O_i = running output accumulator ``` This sample provides: 1. **A practical rocWMMA attention example** -- Shows how to express attention as two fused GEMMs: `Q * K^T` followed by `P * V`. 2. **An educational Flash Attention v2 reference** -- Demonstrates online softmax, Q-tile outer-loop scheduling, causal masking, and register-resident output accumulation. 3. **A fused B2B-GEMM pattern** -- Demonstrates how rocWMMA fragments, LDS staging, cooperative loads, and MFMA accumulators can be composed into a multi-stage attention pipeline. 4. **A modern LLM kernel pattern** -- Covers core ideas used by production attention kernels while keeping the implementation readable as a community sample. **References:** - Dao et al., "FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness", NeurIPS 2022 -- https://arxiv.org/abs/2205.14135 - Dao, "FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning", 2023 -- https://arxiv.org/abs/2307.08691 - Dao-AILab FlashAttention reference implementation -- https://github.com/Dao-AILab/flash-attention ## Technical Details ### Fusion Technique The kernel fuses the full attention forward path into one GPU kernel: | Optimization | Description | Benefit | | ------------------------------- | ------------------------------------------------------------------------------------ | ------------------------------------------------------------------------- | | **Q-tile outer loop** | Each CTA owns one `[Br x D]` Q tile and iterates over the K/V sequence | Matches Flash Attention v2 work partitioning | | **Online softmax** | Maintains running `m_i` and `l_i` per row instead of materializing full scores | Avoids global `[N x N]` score matrix | | **Dual-GEMM pipeline** | GEMM #1 computes `S = Q * K^T`; GEMM ROCm#2 computes `O += P * V` | Expresses attention through rocWMMA MFMA fragments | | **K/V double buffering** | Two LDS K/V buffers ping-pong across sequence tiles | Overlaps global-memory reads with compute | | **Register-resident output** | `O_i` stays in fp32 accumulator fragments across the K/V loop | Avoids global intermediate output buffers | | **LDS-mediated softmax bridge** | `S` is staged through LDS for row-wise softmax and converted to fp16 `P` for GEMM ROCm#2 | Provides a portable bridge between accumulator layout and matrix_a layout | | **Causal tile skipping** | Causal mode skips K/V tiles strictly to the right of the Q tile | Reduces unnecessary work for autoregressive attention | ### Kernel Data-Flow ```text Global Q,K,V --coop load--> LDS | | Q tile loaded once | K/V tiles double-buffered | +-----+-----+ | | v v GEMM #1: Q * K^T K read as col_major = K^T | frag_S | store -> S_lds | online softmax update m_i, l_i, alpha | cast/write -> P_lds | v GEMM ROCm#2: O += P * V | frag_O persists across K/V loop | final normalize: O = O / l_i | cast fp32 -> fp16 | O global ``` ### Online Softmax For each Q row, the kernel processes K/V tiles and updates: ```cpp m_new = max(m_i, rowmax(S_ij)) alpha = exp(m_i - m_new) P_ij = exp(S_ij - m_new) l_new = alpha * l_i + rowsum(P_ij) O_i = alpha * O_i + P_ij * V_j ``` This keeps the result exact without storing the full attention matrix. When a later K/V tile contains a larger row max, the old output accumulator is rescaled by `alpha`. ### Parallel Decomposition | Item | Value | | -------------- | ----------------------------------- | | Tensor layout | contiguous `[B, H, N, D]` row-major | | Grid | `(ceil(N / Br), H, B)` | | Block | `TBLOCK_X` threads, one-dimensional | | Q tile | `Br = 64` rows | | K/V tile | `Bc = 32` rows | | Head dimension | `Dhead = 64` | | Waves per CTA | `WAVES = 4` | | Rows per wave | `Br / WAVES = 16` | Each wave owns one contiguous 16-row stripe of the Q tile: ```text Br = 64 | wave 0 | rows [ 0, 16) | | wave 1 | rows [16, 32) | | wave 2 | rows [32, 48) | | wave 3 | rows [48, 64) | ``` This partition avoids cross-wave communication for the softmax statistics because each wave owns independent rows. ### LDS Layout The kernel uses one dynamic shared memory pool. | Segment | Type | Shape | Bytes | Purpose | | --------- | ---- | ------------------------- | ---------: | --------------------------------------- | | `Q_lds` | fp16 | `[Br x D]` | 8 KiB | Q tile, loaded once | | `KV_lo` | fp16 | `[Bc x D] K + [Bc x D] V` | 8 KiB | K/V ping buffer | | `KV_hi` | fp16 | `[Bc x D] K + [Bc x D] V` | 8 KiB | K/V pong buffer | | `P_lds` | fp16 | `[Br x Bc]` | 4 KiB | Softmax output for GEMM ROCm#2 | | `SO_lds` | fp32 | `[Br x max(Bc, D)]` | 16 KiB | Aliased scratch for S and O rescale | | **Total** | | | **44 KiB** | Fits within standard dynamic LDS budget | `Bc = 32` is chosen instead of 64 because double-buffering both K and V with `Bc = 64` would require about 64 KiB LDS and may need dynamic shared memory opt-in. With `Bc = 32`, the sample keeps K/V double buffering while staying around 44 KiB. ### Key Technique: LDS-Mediated B2B-GEMM Bridge The sample uses two GEMMs: ```cpp S = Q * K^T O += P * V ``` The bridge between them is: 1. Store fp32 accumulator fragment `S` to LDS. 2. Apply scale, causal mask, row max, exp, and row sum in LDS. 3. Cast the softmax probability `P` to fp16. 4. Load `P` as a `matrix_a` fragment for the second GEMM. This mirrors the general B2B-GEMM staging pattern used in fused kernels. It is intentionally portable and readable, even though a production kernel would typically replace some LDS round-trips with cross-lane shuffles. ### rocWMMA Features Demonstrated - **Fragment types** -- `matrix_a`, `matrix_b`, and `accumulator` - **Mixed precision** -- fp16 input/output with fp32 compute - **Cooperative global reads** -- `fragment_scheduler::coop_row_major_2d` - **K/V double buffering** -- ping-pong LDS buffers for sequence tiles - **B2B-GEMM fusion** -- `QK^T -> softmax -> PV` - **Accumulator-resident output** -- keeps `O` in fp32 fragments across K/V loop - **LDS-mediated softmax** -- readable row-wise softmax implementation - **Causal masking** -- scalar mask application during softmax - **Matrix layout trick** -- K is stored row-major but read as `col_major matrix_b`, giving `K^T` without an explicit transpose - **fp32 -> fp16 output cast** -- `MfmaFragAcc` to `MfmaFragStoreOut` ### Supported Architectures | Architecture | Wave Size | TBLOCK_X | Status | | ---------------- | -------------------: | -------: | --------------------- | | gfx1201 / RDNA 4 | wave32 | 128 | Tested on RX 9070 | | gfx942 / CDNA 3 | wave64 expected path | 256 | Intended support path | The sample includes architecture-specific wave-size checks. The tested hardware for this PR is RX 9070 / gfx1201. ### Files Changed | File | Change | | -------------------------------------------- | ----------------------------------------------------------------------- | | `samples/community/simple_flash_attn_v2.cpp` | **[NEW]** Flash Attention v2 forward community sample | | `samples/community/CMakeLists.txt` | **[MODIFIED]** Added `add_community_sample(simple_flash_attn_v2 ...)` | | `samples/community/README.md` | **[MODIFIED]** Added Flash Attention v2 entry to community samples list | ### Known Limitations - Head dimension is fixed to `Dhead = 64`. - Sequence length `N` must be a multiple of both `Br = 64` and `Bc = 32`. - Tensor layout is fixed to contiguous `[B, H, N, D]` row-major. - No backward pass. - No dropout. - No ALiBi. - No sliding-window mask. - Causal masking is supported, but applied scalar-wise during the LDS softmax pass. - Softmax row reductions currently go through LDS; production kernels usually use cross-lane shuffle reductions. - Output rescaling also round-trips through LDS for portability. - Performance is not tuned for production use; this is an educational community sample. - LDS usage is about 44 KiB per block. - Input values are damped by `1/8` to keep softmax logits in a stable fp16 range for validation. ### Build & Run ```bash # Navigate to rocWMMA project root cd <repo>/projects/rocwmma # Configure -- disable tests and official samples for fast build # Use Debug build to enable CPU reference validation cmake -B build . \ -DCMAKE_CXX_COMPILER=/opt/rocm/bin/amdclang++ \ -DCMAKE_BUILD_TYPE=Debug \ -DROCWMMA_BUILD_COMMUNITY_SAMPLES=ON \ -DROCWMMA_BUILD_SAMPLES=OFF \ -DROCWMMA_BUILD_TESTS=OFF \ -DGPU_TARGETS="gfx1201" # Build only the community sample target cmake --build build --target simple_flash_attn_v2 # Quick validation ./build/samples/community/simple_flash_attn_v2 # Run all validation cases ./build/samples/community/simple_flash_attn_v2 --all # Optional: skip CPU reference validation ./build/samples/community/simple_flash_attn_v2 --skip-validation ``` > Note: The monorepo root `CMakeLists.txt` does not include rocWMMA as a component. Use `projects/rocwmma/` as the CMake source directory. ## Test Plan The sample includes two quick validation cases plus two larger GPT-style causal attention cases. | Config | B | H | N | D | Causal | Source | | ------------------------ | -: | -: | ---: | -: | ------ | -------------------------- | | Default non-causal | 1 | 2 | 128 | 64 | no | Quick validation | | Default causal | 1 | 2 | 128 | 64 | yes | Quick validation | | GPT-2-small-style causal | 1 | 12 | 1024 | 64 | yes | Larger LLM-like validation | | Longer-context causal | 1 | 12 | 2048 | 64 | yes | Longer sequence validation | Validation is performed by comparing GPU output against a CPU reference implementation in Debug builds. The tested pass criterion is max relative error below `0.001`. ## Test Result The sample builds successfully, runs successfully, and passes CPU-reference validation on all tested configurations on RX 9070 / gfx1201. *Build/run evidence on RX 9070 / gfx1201 is attached in PR comments.* ### RX 9070 / gfx1201 -- Run Output Summary ```text Community Sample: Flash Attention v2 Forward This sample demonstrates: fused attention with online softmax, dual-GEMM register pipeline, cooperative loads + K/V double buffering -- the core of every modern LLM attention kernel. === GPU Hardware Info === Device name : AMD Radeon RX 9070 GCN arch : gfx1201 Compute units : 28 Warp size : 32 Global memory : 16304 MiB Shared mem/blk : 64 KiB Max clock (MHz) : 2120 Memory bw (GB/s): 80.512 ======================== ``` | Config | B | H | N | D | Causal | elapsedMs | TFlops/s | Max Rel Error | Status | | ------------------------ | -: | -: | ---: | -: | ------ | --------: | -------: | ------------: | ------ | | Default non-causal | 1 | 2 | 128 | 64 | no | 0.179049 | 0.234255 | 0.000159566 | PASSED | | Default causal | 1 | 2 | 128 | 64 | yes | 0.192728 | 0.108814 | 0.000162575 | PASSED | | GPT-2-small-style causal | 1 | 12 | 1024 | 64 | yes | 5.12248 | 1.5721 | 0.000162734 | PASSED | | Longer-context causal | 1 | 12 | 2048 | 64 | yes | 9.46681 | 3.40265 | 0.000162734 | PASSED | The following results are from one local sanity run on gfx1201. They are intended to confirm that the sample builds, runs, and validates correctly, rather than to serve as production performance claims. ### Local Run Metadata | Item | Detail | | ---------- | ------------------------ | | Target | `simple_flash_attn_v2` | | GPU target | `gfx1201` | | Build type | Debug | | Run args | `--all` | | Hardware | AMD Radeon RX 9070 | | Validation | CPU reference comparison | | Result | All cases passed | ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Test Result [result.txt](https://github.com/user-attachments/files/29020225/result.txt) --------- Co-authored-by: Bain Syrowik <bsyrowik@amd.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
20 participants
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.
Summary
Resolves the merge conflicts blocking PR ROCm#6534 by merging the latest
ROCm/rocm-libraries:developinto this branch.The branch had diverged: 39 commits ahead of a merge base that
developwas ~5455 commits past. Three files conflicted, all in the FMHA forward example/test:example/ck_tile/01_fmha/example_fmha_fwd.cppexample/ck_tile/01_fmha/fmha_fwd_runner.hpptest/ck_tile/fmha/test_fmha_fwd.cppResolution
All conflicts were the same "both-added" shape: this branch added a
block_maskparameter whiledevelopindependently addedpack_gqa. Both were kept, with two correctness details:.insert(...)fluent arg-parser chain, dropped the statement-terminating;afterblock_maskso the chain continues intopack_gqa.block_mask_str->pack_gqaat every call site to match the runner signature infmha_fwd_runner.hpp(positional args).Formatting
clang-format-18usingprojects/composablekernel/.clang-format.Verification
git diffof the merge result vsorigin/developshows only this branch'sblock_maskadditions;develop'spack_gqapreserved.How to land
Merge this into
meta/fmha-fwd-block-sparsity; that advances the head branch of ROCm#6534 and clears its conflict state againstdevelop.