Repository navigation
[hipblaslt] Equality tuning for BF16 TN (gfx1250) - #10582
Closed
minsukim-amd wants to merge 52 commits into
Closed
minsukim-amd wants to merge 52 commits into
minsukim-amd wants to merge 52 commits into
Conversation
feat(stinkytofu): insert xcnt drains for gfx1250 XNACK-replay hazards JIRA ID: AIHPBLAS-4203 ## Motivation On gfx1250 an unavailable address translation can return XNACK-Retry: the wave keeps issuing while earlier translations are still in flight, then replays the faulting instruction. If a replay source register was overwritten in the meantime, the replayed access re-executes with the wrong address or data. Hardware does not cover every case, so the remaining ones need an explicit `s_wait_xcnt 0` drain. ## What this does Adds `Gfx1250HazardPass`, which walks the final instruction order and drains XCNT where a replay group would otherwise be unsafe: - FLAT must not overwrite a source of its group (a single-instruction group may overwrite its own source); SMEM must not overwrite any group source. - The first atomic after non-atomic memory starts with XCNT == 0; a consecutive atomic run needs no drain. - Drains before forever `s_sleep`, pre-existing `s_prefetch`, and a `s_set_vgpr_msb` that is not adjacent to its VMEM user. - Ordered Global/Buffer/Scratch/Image traffic needs nothing. The pass runs whole-kernel (entry plus callable functions) for `RequiresXCntForVolatileVMEM` arch capability, carries a group across fall-through block boundaries. It also has an optional stderr summary of inserted drains by rule and by whether the drain landed in a loop and/or a matrix block. SMEM and VMEM source overlapping that force a drain are reported as warnings; a multi-dword SMEM load overwriting its own source cannot be repaired by any drain, so the pass reports it and asserts (single-dword SMEM load is fine). ## Testing `gfx1250_xnack_hazard_test.stir` covers the rules and boundaries, `gfx1250_xcnt_drain_sites.stir` the drain-site report; `Gfx1250HazardPassTest` (death test) and `Gfx1250HazardProfileTest` cover the abort and the kernel/helper count split. ctest suite passes. tox on machine and ffm have no unexpected new failures.
…eaders (#10425) ## Summary Nine `switch` statements in `projects/rocalution/src/base/host/host_io.cpp` pick the conversion to apply when a matrix file's value type differs from the in-memory `ValueType`. In all nine, the `complex32` case has a body but no `break`, so control falls into the `complex64` case: ```cpp case rocsparseio_type_complex32: { copy_mixed_arrays(nnz, val[0], (const std::complex<float>*)tmp_val); } case rocsparseio_type_complex64: { copy_mixed_arrays(nnz, val[0], (const std::complex<double>*)tmp_val); break; } ``` The result is a heap buffer over-read plus silent corruption of the values that were just converted correctly. Related: found during the same monorepo sweep as #10413. ## Why it is a bug Three facts combine, all in-tree: 1. **The scratch buffer is sized from the file's type, not the in-memory type.** `host_io.cpp:918-922`: ```cpp size_t sizeof_val_type; status = rocsparseio_type_get_size(file_val_type, &sizeof_val_type); tmp_val = malloc(nnz * sizeof_val_type); ``` and the `switch` is on `file_val_type`, so reaching the `complex32` case means the buffer is `nnz * 8` bytes. 2. **The two widths differ by 2x.** `src/utils/rocsparseio.hpp:415-421`: ```cpp case complex32: return sizeof(float) * 2; // 8 case complex64: return sizeof(double) * 2; // 16 ``` 3. **`copy_mixed_arrays` reads `size` elements of its source type.** `host_io.cpp:656-667`: ```cpp for(size_t i = 0; i < size; ++i) { x[i] = static_cast<X>(y[i]); } ``` So the fallthrough reads `nnz * 16` bytes from an `nnz * 8` byte allocation: an `nnz * 8` byte over-read. It then writes those garbage values over `val[0]`, discarding the correct conversion the `complex32` case had already performed. ## Reachability, stated precisely Exactly one configuration reaches the fallthrough: **a file whose value type is `complex32`, read into `ValueType` `std::complex<double>`.** I checked the other possibilities rather than assuming: | `ValueType` | What happens | Reaches fallthrough | | --- | --- | --- | | `std::complex<float>` | `required_val_type == complex32 == file_val_type`, so `same_val_type` is true and the whole `switch` is skipped | no | | `std::complex<double>` | `complex32` case runs the real `complex<double> <- complex<float>` conversion, then falls through | **yes** | | `int8_t`, `float`, `double` | the `complex32` case hits a `copy_mixed_arrays` specialization whose body is `throw 1;`, so it throws first | no | The throwing specializations are at `host_io.cpp:696-731`. `std::complex<double>` is explicitly instantiated for every one of these readers (`host_io.cpp:4241`, `4303`, `4361`, `4434`, `4504`, and the CSR/MCSR/BCSR equivalents), so the path is compiled into the shipping library. It is the ordinary mixed-precision case of loading a single-precision complex matrix into a double-precision complex one. I want to be straight that this is one specific type pairing rather than "any mismatched complex read". It is still a live memory-safety defect on a supported and instantiated path. ## Scope Affected readers, one site each except `hyb` which has two (the COO value array and the ELL value array): ``` host_io.cpp:1022 read_matrix_csr_rocsparseio switch on file_val_type host_io.cpp:1330 read_matrix_mcsr_rocsparseio switch on file_val_type host_io.cpp:1681 read_matrix_bcsr_rocsparseio switch on file_val_type host_io.cpp:1967 read_matrix_coo_rocsparseio switch on file_val_type host_io.cpp:2231 read_matrix_dia_rocsparseio switch on file_val_type host_io.cpp:2488 read_matrix_ell_rocsparseio switch on file_val_type host_io.cpp:2858 read_matrix_hyb_rocsparseio switch on file_coo_val_type host_io.cpp:2931 read_matrix_hyb_rocsparseio switch on file_ell_val_type host_io.cpp:3123 read_matrix_dense_rocsparseio switch on file_val_type ``` All nine switches have identical shape. Every case group ends in `break` except `complex32`: ``` switch@996 on file_val_type: break case int32+int64 break case int8 break case float32 break case float64 NO-EXIT case complex32 break case complex64 ``` **`complex32` also appears 13 more times as a stacked case label sharing a body with its neighbours**, for example `host_io.cpp:955-962`: ```cpp case rocsparseio_type_int8: case rocsparseio_type_float32: case rocsparseio_type_float64: case rocsparseio_type_complex32: case rocsparseio_type_complex64: { break; } ``` Those are correct and are **not** touched. Only the nine cases that carry their own body are changed. A first pass of my own scan counted those 13 as defects before I distinguished stacked labels from bodies, so the number in this PR is 9, not 22. ## Fix `break;` added to the nine `complex32` bodies. Nine lines, no other change: ```diff case rocsparseio_type_complex32: { copy_mixed_arrays(nnz, val[0], (const std::complex<float>*)tmp_val); + break; } ``` ## Verification No AMD GPU was used, and none is needed: this is host-side matrix-file I/O. I built a reduced standalone reproducer modelling the buffer sizing, the 8 versus 16 byte type widths, the `copy_mixed_arrays` overload set including the throwing specializations, and the `switch` itself, on the reachable `ValueType = std::complex<double>` instantiation. `nnz = 4`, so `tmp_val` is 32 bytes and the fallthrough reads 64. **Before, `clang++ -std=c++17 -g -O0 -fsanitize=address`:** ``` tmp_val = malloc(4 * 8) = 32 bytes after complex32 case: val[0] = (1,10) (2,11) (3,12) (4,13) <-- correct ================================================================= ==32738==ERROR: AddressSanitizer: heap-buffer-overflow on address 0x603000001c50 at pc 0x000103082c84 bp 0x00016d5d6120 sp 0x00016d5d58d0 READ of size 16 at 0x603000001c50 thread T0 #0 0x000103082c80 in __asan_memcpy+0x400 (libclang_rt.asan_osx_dynamic.dylib:arm64e+0x3ec80) #1 0x0001028290f0 in void copy_mixed_arrays<std::__1::complex<double>, std::__1::complex<double>>(unsigned long, std::__1::complex<double>*, std::__1::complex<double> const*) repro.cpp:53 #2 0x000102828b58 in main repro.cpp:129 #3 0x00018c759d50 in start+0x1c0c (dyld:arm64e+0x8d50) 0x603000001c50 is located 0 bytes after 32-byte region [0x603000001c30,0x603000001c50) allocated by thread T0 here: #0 0x000103085164 in malloc+0x78 (libclang_rt.asan_osx_dynamic.dylib:arm64e+0x41164) #1 0x000102828980 in main repro.cpp:82 #2 0x00018c759d50 in start+0x1c0c (dyld:arm64e+0x8d50) ``` Note the frame: the over-read is attributed to the `copy_mixed_arrays<complex<double>, complex<double>>` instantiation, which is the generic template reached only through the fallthrough. At `-O1` the loop vectorises and ASan reports the same overflow as a single `READ of size 64`. **Before, no sanitizer**, showing the second half of the defect: ``` tmp_val = malloc(4 * 8) = 32 bytes after complex32 case: val[0] = (1,10) (2,11) (3,12) (4,13) <-- correct after complex64 case: val[0] = (524288,1.04858e+06) (2.09715e+06,4.19431e+06) (0,0) (0,0) <-- clobbered done ``` The correct values are replaced by the bit patterns of pairs of `std::complex<float>` reinterpreted as single `std::complex<double>` values, trailed by the out-of-bounds tail. **After, `-DFIXED`, ASan clean at both `-O0` and `-O1`:** ``` tmp_val = malloc(4 * 8) = 32 bytes after complex32 case: val[0] = (1,10) (2,11) (3,12) (4,13) <-- correct done EXIT=0 ``` ### Why no compiler caught it `-Wall -Wextra` are silent. `-Wimplicit-fallthrough` does flag it, but clang does not imply it from either, and rocALUTION's CMake never enables it. Turning it on for `projects/rocalution` would prevent recurrence, and I am happy to do that as a follow-up if you want it separated from this fix. ## A note on the unit test policy check `tools/libraries_pr_bot/policy.yml` requires an accompanying test file for `.cpp` changes, with `exempt_paths: []`. I have not added one, and I would rather explain why than route around the check. `projects/rocalution/clients/tests/` has 31 gtest files and **none of them cover file I/O at all**: `git grep -l "ReadFileRSIO\|WriteFileRSIO\|rocsparseio" -- projects/rocalution/clients/` returns nothing. So there is no existing test file to extend, and a new one needs a rocsparseio fixture written at a complex32 value type, a new entry in `clients/tests/CMakeLists.txt`, and `init_rocalution()`. I cannot build or run the rocALUTION client suite here, and shipping unexercised CMake and gtest code alongside a memory-safety fix would make this PR harder to trust, not easier. Two things I am happy to add, whichever you prefer: 1. A `WriteFileRSIO` then `ReadFileRSIO` round-trip test in a new `clients/tests/test_host_io.cpp`, writing a small `complex<float>` matrix and reading it back into a `LocalMatrix<std::complex<double>>`, asserting the values match. That is the faithful regression test and it fails on `develop` under ASan. I would need someone to run it once, since I have no ROCm runtime. 2. Enabling `-Wimplicit-fallthrough` for `projects/rocalution`, which pins the whole class rather than this instance. I can also hand over the standalone reproducer if it is useful for the review. ## Checklist - Based on `develop` (`86a82b0`), single commit, DCO signed off. - Re-verified against current `develop` HEAD rather than an earlier snapshot: 9 affected sites, 13 stacked-label sites correctly excluded, all other cases in all nine switches confirmed to still end in `break`. - Prior art checked: no open PR or issue mentions `host_io`, `copy_mixed_arrays` or `rocsparseio_type_complex32`, and `git log -S 'rocsparseio_type_complex32'` shows no partial fix. - `projects/rocalution` is **not** in the root `.pre-commit-config.yaml` exclude list, so `clang-format` 18.1.4, the pinned version, was run against `projects/rocalution/.clang-format`. It reformats nothing beyond the nine added lines. No trailing whitespace, file still ends in a newline. - Behaviour is unchanged for every configuration except the one that was reading out of bounds. Signed-off-by: Aditya Singh <adisin650@gmail.com>
…alution/docs/sphinx (#10477) Bumps [cryptography](https://github.com/pyca/cryptography) from 48.0.1 to 50.0.0. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/pyca/cryptography/blob/main/CHANGELOG.rst">cryptography's changelog</a>.</em></p> <blockquote> <p>50.0.0 - 2026-07-31</p> <pre><code> * **SECURITY ISSUE**: :func:`~cryptography.hazmat.primitives.serialization.pkcs7.pkcs7_decrypt_der` and its PEM and S/MIME variants no longer expose distinguishable errors or timing when unwrapping a ``RecipientInfo``'s ``encryptedKey``, which could act as a Bleichenbacher oracle for callers that decrypt untrusted messages. A random key is now substituted on failure, as described in :rfc:`3218`. Credit to **@X1AOxiang** for reporting the issue. **CVE-2026-69247** * Deprecated Diffie-Hellman key exchange over finite fields (FFDH). Everything FFDH is deprecated, including the types in ``cryptography.hazmat.primitives.asymmetric.dh`` and loading FFDH keys or parameters with the key loading APIs. Users should migrate to a more modern key exchange algorithm. * Added ``xof()`` class methods to :class:`~cryptography.hazmat.primitives.hashes.SHAKE128` and :class:`~cryptography.hazmat.primitives.hashes.SHAKE256` for constructing algorithm instances configured for use with :class:`~cryptography.hazmat.primitives.hashes.XOFHash`. * The :mod:`X.509 verification <cryptography.x509.verification>` APIs are now considered stable and are subject to our API stability policy. * Added the :doc:`/cobblestone` recipe, an implementation of the Cobblestone-128 and Cobblestone-256 instantiations of the `C2SP chunked-encryption specification <https://c2sp.org/chunked-encryption>`_ for streaming authenticated encryption of large messages. * Parsing a Signed Certificate Timestamp list now rejects encodings that carry trailing bytes after the list or after an individual SCT, instead of silently ignoring them. * Added support for using :class:`~cryptography.x509.Name` as a field type in the :doc:`/hazmat/asn1/index` module. * Loading a public key or an EC private key now rejects DER where the ``subjectPublicKey`` (or EC ``publicKey``) ``BIT STRING`` declares a non-zero number of unused bits, instead of silently ignoring it. * Parsing a CRL entry's ``InvalidityDate`` extension now rejects a ``GeneralizedTime`` that carries fractional seconds or another non-DER form, matching the strict encoding already required for every other X.509 time field. * :func:`~cryptography.x509.ocsp.load_der_ocsp_request` and :func:`~cryptography.x509.ocsp.load_der_ocsp_response` now reject a request or response whose ``version`` field is not ``v1``, the only version defined by RFC 6960, matching the version validation already performed when loading certificates, CSRs and CRLs. * :class:`~cryptography.hazmat.primitives.hashes.XOFHash` is now supported when building against AWS-LC. * HMAC (and therefore PBKDF2-HMAC) with SHA-3 hashes is now supported when building against AWS-LC. * Diffie-Hellman (:doc:`/hazmat/primitives/asymmetric/dh`) is now supported when building against AWS-LC. </tr></table> </code></pre> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/pyca/cryptography/commit/dcb7050b807b00392fa9fe2eac7cb362fcf355cc"><code>dcb7050</code></a> Prepare for 50.0.0 release (<a href="https://github.com/pyca/cryptography/issues/15372">#15372</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/53fccd93413a8d7f07d6d8999681f27b75cffa3f"><code>53fccd9</code></a> Don't leak how PKCS#7 encryptedKey decryption failed (<a href="https://github.com/pyca/cryptography/issues/15369">#15369</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/d472f978470fbefa521b86d98b2ecccbbb4d1dd8"><code>d472f97</code></a> Add <code>from __future__ import annotations</code> to all src/ Python files (<a href="https://github.com/pyca/cryptography/issues/15371">#15371</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/908773d53829fb1466c6db364b31321c3cd8eb9a"><code>908773d</code></a> Bump downstream dependencies in CI (<a href="https://github.com/pyca/cryptography/issues/15368">#15368</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/2cc07cc948948211899bcb0cddd1fddf86e95812"><code>2cc07cc</code></a> Bump BoringSSL, OpenSSL, AWS-LC in CI (<a href="https://github.com/pyca/cryptography/issues/15367">#15367</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/c94ede9f040fa44942f7139772603419000acf66"><code>c94ede9</code></a> chore(deps): bump ruff from 0.16.0 to 0.16.1 (<a href="https://github.com/pyca/cryptography/issues/15366">#15366</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/67a8308dc9ea4cce6056e0f1438f903c208c3f35"><code>67a8308</code></a> chore(deps): bump virtualenv from 21.7.0 to 21.7.1 (<a href="https://github.com/pyca/cryptography/issues/15365">#15365</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/95018ffcdbbc510fd92fc872e3a3e80aa6e58596"><code>95018ff</code></a> Release the GIL in one-shot AEAD encrypt/decrypt (<a href="https://github.com/pyca/cryptography/issues/15361">#15361</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/6954733eaf55a0074abf88f06f7242dfca3a5d02"><code>6954733</code></a> Release the GIL during DH and DSA parameter generation (<a href="https://github.com/pyca/cryptography/issues/15364">#15364</a>)</li> <li><a href="https://github.com/pyca/cryptography/commit/6893b94c33e948f6240082461424cfb5da2dacc6"><code>6893b94</code></a> Import _serialization instead of serialization in x509/extensions (<a href="https://github.com/pyca/cryptography/issues/15363">#15363</a>)</li> <li>Additional commits viewable in <a href="https://github.com/pyca/cryptography/compare/48.0.1...50.0.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 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 The docs build aborted with a fatal Doxygen error: `rocsparse_spmat_scale.h:57: error: unable to resolve reference to 'rocsparse_dnvec_descr_create_scalar' for \ref command (warning treated as error, aborting now)` ## Technical details `rocsparse_spmat_scale.h` unconditionally documents `rocsparse_spmat_scale` and uses `\ref rocsparse_dnvec_descr_create_scalar`, but that referenced function is declared inside `#ifdef ROCSPARSE_WITH_SPMAT_SCALE` in `rocsparse-auxiliary.h`. Since the Doxygen build did not have `ROCSPARSE_WITH_SPMAT_SCALE` defined, the referenced symbol was invisible to Doxygen and the `\ref` could not resolve — which is fatal with warnings-as-errors. Added `ROCSPARSE_WITH_SPMAT_SCALE` to the `PREDEFINED` list in `docs/doxygen/Doxyfile`, consistent with how the sibling feature flags `ROCSPARSE_WITH_HANDLE_CREATE` and `ROCSPARSE_WITH_ILDLT0` are already handled. This makes Doxygen parse the guarded declaration so the cross-reference resolves and the feature documents consistently. JIRA ID: AISPARSE-739
## Motivation Somehow the old algorithm used a over complicated logic to count the buckets, after the exclusive scan, each thread already knows its bucket offset and size, so it can directly check whether rank falls into its bucket range. This MR removes the extra inclusive scan and shared offset array without changing the selection logic. JIRA ID: EXSWSTRHPC-276 ## Test Result <img alt="nth_90a" src="https://github.com/user-attachments/assets/260d9881-b1b3-4df0-a0b6-7ead51c83c7d" /> <img alt="nth_942" src="https://github.com/user-attachments/assets/914b9bdd-a622-4723-b247-d0c4c14656da" /> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Summary JIRA ID : ALMIOPEN-2041 Feature Flag Protected: Yes - HIPDNN_ENABLE_CUDNN_COMPATIBILITY Completes the cuDNN-frontend graph node surface on the compatibility shim added in #9361, so hipified cuDNN v9 consumers can build the full set of `graph::Graph::*` node methods against hipDNN. Tier-1 nodes with a real hipDNN engine forward to the wrapped frontend graph; the remaining cuDNN v9 nodes are fail-stubs that record `GRAPH_NOT_SUPPORTED` (surfaced at the next `validate()`) while still handing back a live, graph-registered placeholder tensor — so hipified source compiles, survives the idiomatic `node(...)->set_output(true).set_uid(n)` chain, and fails loudly at `validate()` instead of dereferencing null. The shim targets cuDNN frontend **v1.24.0** (`CUDNN_FRONTEND_VERSION == 12400`, declared in `hipdnn_compatibility/cudnn/cudnn_frontend_version.h`). `detail/graph_wrapper.h` `static_assert`s on that macro, so the pin cannot move without a deliberate re-diff of every node signature against upstream `graph_interface.h` / `node_interface.h`. ## Risk Assessment Risk 3 (low-moderate). Header-only and host-only: no kernel selection, dispatch, or shipping default changes. The cuDNN-shaped surface — new public header, ~18 Tier-1 node methods, 22 fail-stub attribute classes — is behind the opt-in `HIPDNN_ENABLE_CUDNN_COMPATIBILITY` flag (default OFF). Two cuDNN spellings land **unguarded** in the native `hipdnn_frontend::graph` namespace regardless of the flag: `Layernorm_backward_attributes` and `Resample_attributes`, `typedef`s of existing types added alongside the already-unguarded `Resample_fwd_attributes`. That is intentional — cuDNN spellings are welcome in the native namespace — not a side effect of the shim. Tier-1 node methods are thin forwarders to existing, already-tested hipDNN engines; the only new runtime behavior is the fail-stub record-and-return-placeholder path, which is host-covered. Also changes the superbuild CI workflow (enables the flag in both jobs). ## ASIC Coverage ASIC-independent. Frontend source-compatibility plumbing that lowers cuDNN-spelled graphs onto the existing hipDNN frontend; changes no kernel selection, support surface, or default behavior. All tests validate host-side with no device; the four samples are host-only graph-build+validate. No multi-arch sweep required. The shim, its tests, and its samples build only with `HIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON`, and this PR is what turns that on in CI: it adds `-DHIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON` to both jobs in `hipdnn-superbuild-ci.yml`, both of which run `ctest --output-on-failure`. Shim coverage is therefore exercised by PR CI on gfx1151/Windows and gfx94X/Linux, not just locally. ## Testing Summary - Node-coverage gtests (`TestCudnnShimGraphNodes.cpp`): Tier-1 nodes build a well-shaped graph that `validate()`s good; Tier-2 fail-stubs record `GRAPH_NOT_SUPPORTED` with the issue-tracker message; a poison test proves a fail-stub overrides an otherwise-valid graph (first-error-wins); a compile-time check that all 39 cuDNN v9 `*_attributes` classes construct and chain `.set_name`. - Fail-stub return-value tests: a single-output stub (`reshape`) and multi-output stubs (`genstats`, `sdpa_fp8`) hand back live placeholders that survive `->set_output(true).set_uid(n)` before `validate()` reports `GRAPH_NOT_SUPPORTED` — the shape that a null return would crash on. - `sdpa_fp8_backward` overload tests: both the 18-tensor FP8 form (7 outputs) and the 17-tensor MXFP8 form (6 outputs) resolve and record, pinning arity and output count against the FE v1.24.0 signatures. - Host-only samples registered as CTests, mirroring upstream cuDNN FE sample names: ConvFprop, Matmul, Layernorm, Pointwise. ## Testing Checklist - [x] hipDNN unit tests (cuDNN shim node coverage) - built with `HIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON`, ran `hipdnn-unit-check` - Status: Passed - [x] cuDNN-shim host-only samples - `ctest -R hipdnn_sample_cudnn_shim_` - Status: Passed - [x] PR CI (builds and runs the shim tests and samples with the flag ON) - GitHub PR checks - Status: Passed ## Technical Changes - Adds the Tier-1 node methods to `detail/graph_wrapper.h` (conv fprop/dgrad/wgrad, batchnorm + backward + inference, layernorm + backward, rmsnorm + backward, matmul, pointwise, reduction, resample, block-scale quantize/dequantize), each forwarding to the wrapped hipDNN graph and flipping the graph to Native mode. - Adds `detail/node_wrappers/unsupported_nodes.h`: a CRTP `UnsupportedAttributes` base plus macros that stamp Tier-2 fail-stub attribute classes and matching node methods recording `GRAPH_NOT_SUPPORTED` for the cuDNN v9 nodes with no hipDNN engine yet. Fail-stubs return graph-registered placeholder tensors (single or array, per the node's upstream return type) so the recorded error surfaces at `validate()` rather than as a null dereference at the call site. - Pins the targeted cuDNN FE version: `detail/graph_wrapper.h` `static_assert`s on `CUDNN_FRONTEND_VERSION` (v1.24.0), making node-signature drift against upstream a build failure rather than a source-compatibility break for consumers. - Corrects `sdpa_fp8_backward` to upstream's two overloads — 18 tensor params returning `array<...,7>` (FP8) and 17 returning `array<...,6>` (MXFP8) — matching the existing `sdpa_fp8` pair. - Aliases the Tier-1 attribute types into `<shim>::graph` in `cudnn_frontend/graph_properties.h`. - Adds cuDNN-spelled aliases `Layernorm_backward_attributes` and `Resample_attributes` on the native attribute headers (unguarded, as with `Resample_fwd_attributes`). - Enables `HIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON` in both jobs of `.github/workflows/hipdnn-superbuild-ci.yml`, so PR CI builds and runs the shim tests and samples. - Wires the node-coverage gtest and four host-only samples behind `HIPDNN_ENABLE_CUDNN_COMPATIBILITY` in the frontend-tests and samples CMake. - Moves the `_mode = Mode::Native;` flip after the forwarded call in all 18 Tier-1 node methods, matching the pre-existing `sdpa`/`sdpa_backward` ordering: a throwing forward can no longer leave the wrapper claiming Native with no node added. - Adds `get_compute_data_type()` to the Tier-2 `UnsupportedAttributes` base, so a stub exposes the same universal accessor pair as the hipDNN attribute types the Tier-1 aliases resolve to. - Guards all four new samples with the same `try`/`catch` as `SdpaRoundTrip.cpp`.
## Motivation Update 44 underperforming sizes for gfx942 80cu grid. F32/TN. Ticket: ROCM-26088 Co-authored-by: Mark Kim <markkim@amd.com>
…er mode (#10499) ## Summary Stops installing the three rocKE Python CTest entries (`rocke_python_import_smoke`, `rocke_golden_static`, `rocke_pytest`) in provider mode, unblocking every `hipkernelprovider` PR in this repo. These entries currently fail on both Linux and Windows with `ModuleNotFoundError: No module named 'rocke'`. This is a temporary CI unblock, not a fix: it is reverted once [ROCm/TheRock#7145](ROCm/TheRock#7145) lands and the wheel-staging order bug is fixed. ISSUE ID : #10497 ## Risk Assessment Low risk. CI-only, scoped to the provider-mode rocKE install and confined to `if(HIPKERNELPROVIDER_ENABLE_ROCKE)`. It removes currently-100%-failing test entries, so it cannot mask a regression that was passing. No product code, public API, build wiring, or submodule movement. The residual risk is coverage loss: the rocKE Python suite stops running here until reverted, which is why the revert condition is spelled out in the code comment and the commit message. ## ASIC Coverage Passing PR CI is sufficient. The change is ASIC-independent: it gates a set of `install()` calls and pre-registered CTest entries, and touches neither kernel selection, support surface, nor default runtime behavior. The removed tests are host-only Python checks with no GPU dependency, and the flag guarding them is arch-agnostic. ## Testing Summary - Configured `rocke/platform` in provider mode both ways and diffed the generated installed `CTestTestfile.cmake.install`, confirming exactly the three failing entries disappear and all five C++ rocKE smoke tests remain. - Diffed the generated `cmake_install.cmake` to confirm the Python payload (smoke script, golden test, parity harness, golden JSON, `tests/library` packages, pytest tree) is no longer staged. - Verified the change propagates as a normal variable across `add_subdirectory` (CMP0077 NEW under this project's 3.25.2 floor) via a wrapper project reproducing the real call shape. - Verified the categories YAML needs no edit: the generated labeling code wraps every test in `if(TEST ...)`, so the three now-absent names are a no-op. ## Testing Checklist - [x] Generated CTest entries, provider mode - `cmake -S platform -B ... -DROCKE_INSTALL_PROVIDER_CTEST_SUBDIR=hip_kernel_provider` with and without `-DROCKE_INSTALL_PYTHON_TESTS=OFF` - 8 entries before, 5 after; only the 3 failing Python entries removed - Status: Passed - [x] Install rules for the Python payload - grep of generated `cmake_install.cmake` - 5 payload rules before, 0 after - Status: Passed - [x] Normal-variable override across `add_subdirectory` - wrapper project mirroring `dnn-providers/hip-kernel-provider/CMakeLists.txt` - 0 Python entries, 0 payload rules - Status: Passed - [x] Category YAML tolerance for absent tests - `python3 shared/ctest/parse_ctest_categories.py ROCKE_ENGINE_test_categories_external.yaml <install-file>` - emits `if(TEST ${_test})` guards - Status: Passed - [x] pre-commit - `pre-commit run --files dnn-providers/hip-kernel-provider/CMakeLists.txt` - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - `dnn-providers/hip-kernel-provider/CMakeLists.txt`: sets `ROCKE_INSTALL_PYTHON_TESTS OFF` inside the `if(HIPKERNELPROVIDER_ENABLE_ROCKE)` block, alongside the existing `ROCKE_INSTALL_PROVIDER_CTEST_SUBDIR` and `ROCKE_INSTALL_STANDALONE_CTEST` provider-mode settings. - This is the option's already-documented provider-mode behavior. `rocke/platform/CMakeLists.txt` declares it with the comment "Provider mode turns them OFF until their test-artifact packaging is wired up", but nothing in the tree was actually setting it, so it stayed `ON` and the entries were registered against an artifact that no longer carries the `rocke` package. - Standalone rocKE builds are unaffected and keep the Python tests `ON`; the option default is unchanged. - The five C++ rocKE smoke tests (`rocke_ir_serialize_roundtrip`, `rocke_future_intrinsic_lowering`, `rocke_flavor_datalayout_ssot`, `rocke_mma_frag_ssot`, `rocke_tiled_attention_2d_reentrancy`) still build, install, and run. - `ROCKE_INSTALL_PYTHON_PACKAGE` is intentionally left `ON`. It is a separate gate, the install is harmless, and leaving it keeps the diff to the minimum needed to unblock CI. - The in-code comment records both root causes and the exact revert condition so this does not become permanent by accident. ### Why this is driven from rocm-libraries `HIPKERNELPROVIDER_ENABLE_ROCKE` defaults `OFF` in TheRock's `FLAGS.cmake`, so TheRock's own PR CI never runs these entries and stays green. This repo sets `-DTHEROCK_FLAG_HIPKERNELPROVIDER_ENABLE_ROCKE=ON` unconditionally for the component, so the entire flag-gated breakage lands on `hipkernelprovider` PRs here. The CTest registration is emitted by this repo's CMake, which makes it the only side that can suppress the entries without waiting on a TheRock merge and a merge-base ref advance. Full root-cause analysis, including the wheel-staging order bug in `rocke/CMakeLists.txt`, is in [issue #10497](#10497).
…#9785) ## Summary Adds OCP MX FP4 (E2M1) input support to the hipBLASLt provider's `dequantize → GEMM` path, including mixed FP4 + FP8-OCP A/B operands, so hipDNN consumers can run MXFP4 block-scaled GEMMs through the standard graph API. JIRA ID : ALMIOPEN-1800 ## Risk Assessment Medium risk. - The MX FP4 feature is gated behind `HIPBLASLTPROVIDER_ENABLE_MX_GEMM` (default OFF), so it does not change default behavior; the existing FP8 OCP path is unchanged. - This PR also carries the #9962 workaround, which early-returns `false` from `HipblasltMatmulPlanBuilder::isApplicable` on **Windows gfx115x**. In a default (flag-OFF) Windows build on gfx115x, the hipBLASLt *plain* matmul graph now reports unsupported. ## Related - Workaround for: #9962 (hipBLASLt gfx115x Windows GEMM crash) - TheRock flag declaration: ROCm/TheRock#6966 ## ASIC Coverage MX FP4 (E2M1) support is arch-scoped to gfx950/gfx12. Verified on gfx950. ## Testing Summary - Provider unit tests: FP4 type conversion, `isTypeMxOcp`, and MX plan-builder `isApplicable`/`buildPlan` for FP4 and mixed A/B. - MX matmul GPU integration: FP4→{F16,BF16,F32}, mixed FP8/FP4, and FP8-regression, each validated against an independent CPU reference within `getMxTolerance`. - Test-SDK/CPU unit tests: packed FP4 tensor and sub-byte tensor creation. ## Testing Checklist - [x] Commit hooks - `git commit` - [x] Provider unit tests - `./build/bin/hipblaslt_plugin_tests` - ASICs: gfx950 - [x] MX matmul integration - `./build/bin/hipdnn_integration_tests --gtest_filter='*MxMatmul*'` - ASICs: gfx950 - [x] CPU SDK/unit tests - `hipdnn_data_sdk_tests` / `hipdnn_test_sdk_tests` / `hipdnn_frontend_tests` ## Technical Changes - `HipblasltUtils`: map `FP4_E2M1 → HIP_R_4F_E2M1`; add `isTypeMxOcp()` (FP8 E4M3/E5M2 + FP4 E2M1). - `HipblasltMxMatmulPlanBuilder::checkImplementationLimitations`: accept MX OCP inputs with A/B evaluated independently (mixed FP4+FP8 allowed). - New `PackedFp4Tensor` (2 values/byte device layout) + `createTensor(..., packSubByteElements)`; integration harness gives the GPU bundle a packed FP4 buffer and the CPU-reference bundle an unpacked one from the same seed. - MX matmul fixtures collapsed: `MxMatmul` now derives from `MxMatmulMixed`. - Docs: FP4 inputs + mixed A/B in provider and hipDNN `OperationSupport.md`. --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
…ipsparselt/docs/sphinx in the hipsparselt-docs-dependencies group across 1 directory (#8636) Bumps the hipsparselt-docs-dependencies group with 1 update in the /projects/hipsparselt/docs/sphinx directory: [rocm-docs-core](https://github.com/ROCm/rocm-docs-core). Updates `rocm-docs-core` from 1.32.1 to 1.38.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.38.0 (2026-07-16)</h2> <h3>Feat</h3> <ul> <li><strong>theme</strong>: update rocm theme with new secondary nav</li> <li><strong>theme</strong>: add ai-ecosystem theme with new secondary header nav</li> </ul> <h3>Fix</h3> <ul> <li><strong>theme</strong>: update instinct header</li> <li>fix linting</li> <li>expand if statement to include ROCM_LATEST_RELEASE case</li> <li>delete unused variables to fix linting</li> <li>remove preview banner</li> <li><strong>css</strong>: remove unnecessary z-index override on secondary sidebar</li> <li>pin sphinx-book-theme to 1.1.4</li> <li>update nav links</li> <li><strong>theme</strong>: update left side menu title</li> <li><strong>css</strong>: sidebar collapser and breadcrumbs alignment</li> <li>clean up header nav</li> <li><strong>css</strong>: article header icon alignment</li> </ul> <p>[main 41eef9e] bump: version 1.37.0 → 1.38.0 3 files changed, 26 insertions(+), 4 deletions(-)</p> <h2>v1.37.0 (2026-07-13)</h2> <h3>Feat</h3> <ul> <li><strong>theme.py</strong>: surpress sidebar for rocm-ft flavor</li> <li>remove left side bar</li> <li>remove TOC side bar</li> <li>add new flavor for Fluid Topics docs</li> </ul> <p>[main 409e250] bump: version 1.36.0 → 1.37.0 3 files changed, 13 insertions(+), 4 deletions(-)</p> <h2>v1.36.0 (2026-07-08)</h2> <h3>Feat</h3> <ul> <li><strong>llms</strong>: convert raw HTML tables to Markdown in the full text</li> <li><strong>llms</strong>: note that each project publishes its own llms files</li> <li><strong>llms</strong>: add rocm_docs_llms_full_exclude to skip pages from full text</li> <li><strong>llms</strong>: enable generation for rocm-docs-core's own docs</li> <li><strong>llms</strong>: generate llms.txt and llms-full.txt from the resolved doctree</li> <li><strong>theme</strong>: add instinct-design flavor with AMD Instinct Design System</li> </ul> <h3>Fix</h3> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </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.38.0 (2026-07-16)</h2> <h3>Feat</h3> <ul> <li><strong>theme</strong>: update rocm theme with new secondary nav</li> <li><strong>theme</strong>: add ai-ecosystem theme with new secondary header nav</li> </ul> <h3>Fix</h3> <ul> <li><strong>theme</strong>: update instinct header</li> <li>fix linting</li> <li>expand if statement to include ROCM_LATEST_RELEASE case</li> <li>delete unused variables to fix linting</li> <li>remove preview banner</li> <li><strong>css</strong>: remove unnecessary z-index override on secondary sidebar</li> <li>pin sphinx-book-theme to 1.1.4</li> <li>update nav links</li> <li><strong>theme</strong>: update left side menu title</li> <li><strong>css</strong>: sidebar collapser and breadcrumbs alignment</li> <li>clean up header nav</li> <li><strong>css</strong>: article header icon alignment</li> </ul> <h2>v1.37.0 (2026-07-13)</h2> <h3>Feat</h3> <ul> <li><strong>theme.py</strong>: surpress sidebar for rocm-ft flavor</li> <li>remove left side bar</li> <li>remove TOC side bar</li> <li>add new flavor for Fluid Topics docs</li> </ul> <h2>v1.36.0 (2026-07-08)</h2> <h3>Feat</h3> <ul> <li><strong>llms</strong>: convert raw HTML tables to Markdown in the full text</li> <li><strong>llms</strong>: note that each project publishes its own llms files</li> <li><strong>llms</strong>: add rocm_docs_llms_full_exclude to skip pages from full text</li> <li><strong>llms</strong>: enable generation for rocm-docs-core's own docs</li> <li><strong>llms</strong>: generate llms.txt and llms-full.txt from the resolved doctree</li> <li><strong>theme</strong>: add instinct-design flavor with AMD Instinct Design System</li> </ul> <h3>Fix</h3> <ul> <li><strong>llms</strong>: preserve <!-- raw HTML omitted --> line breaks in HTML table cells and format</li> <li><strong>llms</strong>: HTML table conversion fix</li> <li><strong>llms</strong>: de-duplicate TOC pages on the normalized docname</li> <li><strong>llms</strong>: only generate llms files for HTML builders</li> <li><strong>llms</strong>: index generated/doxygen pages instead of dropping them</li> <li><strong>llms</strong>: log unsupported-node warnings as INFO, not WARNING</li> </ul> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/41eef9ed26d57ee5b4de2a18ded61f8c5780f316"><code>41eef9e</code></a> bump: version 1.37.0 → 1.38.0</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/569a6ad2db2fe6baae0381daee5caceb4572c303"><code>569a6ad</code></a> Merge pull request <a href="https://github.com/ROCm/rocm-docs-core/issues/1575">#1575</a> from ROCm/ai-ecosystem-theme</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/4c9d1ae5ee8c1351ffa5eee07f4276171a9f1330"><code>4c9d1ae</code></a> fix(theme): update instinct header</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/94e12b3f06d90c66e600e193d8aa8c6e97906237"><code>94e12b3</code></a> fix: fix linting</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/cc8bcafd37ffd0219c4f7cdf9b04a60993039c84"><code>cc8bcaf</code></a> fix: expand if statement to include ROCM_LATEST_RELEASE case</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/95e8d77ce6dc199b5c434708fbff15f02599d76c"><code>95e8d77</code></a> fix: delete unused variables to fix linting</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/914b92a67801d763f26faabe7c943a8b0962d649"><code>914b92a</code></a> fix: remove preview banner</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/6fba82f4d6fdc026d91ae74e0fc5d71d6b4e72fb"><code>6fba82f</code></a> fix(css): remove unnecessary z-index override on secondary sidebar</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/416cbea2a81451806dcbd6bd123eebe3b2a7af7e"><code>416cbea</code></a> fix: pin sphinx-book-theme to 1.1.4</li> <li><a href="https://github.com/ROCm/rocm-docs-core/commit/43b88b12d35176a7f00e058fc25aeecd4449fbf0"><code>43b88b1</code></a> fix: update nav links</li> <li>Additional commits viewable in <a href="https://github.com/ROCm/rocm-docs-core/compare/v1.32.1...v1.38.0">compare view</a></li> </ul> </details> <br /> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…docs/sphinx (#10523) Bumps [gitpython](https://github.com/gitpython-developers/GitPython) from 3.1.57 to 3.1.58. <details> <summary>Release notes</summary> <p><em>Sourced from <a href="https://github.com/gitpython-developers/GitPython/releases">gitpython's releases</a>.</em></p> <blockquote> <h2>Security and Fixes</h2> <h2>What's Changed</h2> <ul> <li>test: allow Python startup before timeout by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2195">gitpython-developers/GitPython#2195</a></li> <li>Resolve Windows hook Bash through PATH by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2199">gitpython-developers/GitPython#2199</a></li> <li>Revert "Resolve Windows hook Bash through PATH" by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2200">gitpython-developers/GitPython#2200</a></li> <li>test: focus pytest summary on unexpected failures by <a href="https://github.com/raisulchowdhury"><code>@raisulchowdhury</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2203">gitpython-developers/GitPython#2203</a></li> <li>Validate submodule names before filesystem operations by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2202">gitpython-developers/GitPython#2202</a></li> <li>Handle invalid submodule names during recursive updates by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2205">gitpython-developers/GitPython#2205</a></li> <li>Fix Windows path handling on Python 3.13+ by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2196">gitpython-developers/GitPython#2196</a></li> <li>fixup 2202 by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2206">gitpython-developers/GitPython#2206</a></li> <li>Correct Windows hook Bash precedence by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2201">gitpython-developers/GitPython#2201</a></li> <li>Harden config and Git option validation by <a href="https://github.com/Byron"><code>@Byron</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2204">gitpython-developers/GitPython#2204</a></li> <li>Skip tests that need symlink privileges instead of erroring by <a href="https://github.com/Cyrus580529"><code>@Cyrus580529</code></a> in <a href="https://github.com/gitpython-developers/GitPython/pull/2197">gitpython-developers/GitPython#2197</a></li> </ul> <h2>New Contributors</h2> <ul> <li><a href="https://github.com/raisulchowdhury"><code>@raisulchowdhury</code></a> made their first contribution in <a href="https://github.com/gitpython-developers/GitPython/pull/2203">gitpython-developers/GitPython#2203</a></li> <li><a href="https://github.com/Cyrus580529"><code>@Cyrus580529</code></a> made their first contribution in <a href="https://github.com/gitpython-developers/GitPython/pull/2197">gitpython-developers/GitPython#2197</a></li> </ul> <p><strong>Full Changelog</strong>: <a href="https://github.com/gitpython-developers/GitPython/compare/3.1.57...3.1.58">https://github.com/gitpython-developers/GitPython/compare/3.1.57...3.1.58</a></p> </blockquote> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/gitpython-developers/GitPython/commit/30be45d786e95023e23c616fec5cbabca861b44c"><code>30be45d</code></a> prepare changelog for upcoming release</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/fc2f02cc15af0457896f40a8a72b9f1866d75788"><code>fc2f02c</code></a> Merge pull request <a href="https://github.com/gitpython-developers/GitPython/issues/2197">#2197</a> from Cyrus580529/shared-symlink-guard</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/b10e2501438e31cbc1220c4879fc3747c34dccdb"><code>b10e250</code></a> test: use the shared guard instead of local copies</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/e3e5da8476d9e184ebc729d2d85ea035e582d467"><code>e3e5da8</code></a> test: skip tests that need symlink privileges</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/30d05e30c0b7f28bed7e10cb17b10b68e0708af9"><code>30d05e3</code></a> test: add a shared symlink capability guard</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/9a8f6fef579912bc8b10ab513c33ca6ea1f687ef"><code>9a8f6fe</code></a> Merge pull request <a href="https://github.com/gitpython-developers/GitPython/issues/2204">#2204</a> from gitpython-developers/security-fixes</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/f2550b65bf60ca087190981e2c7b6865e201f40c"><code>f2550b6</code></a> Guard pathspec file inputs in high-level commands</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/d9ddb55bdc66ffe8c9932fe460e6b8c8211e47c7"><code>d9ddb55</code></a> Guard unsafe git init options</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/9b5dcaf85da5946dbf69dcd53f9edba08f760b32"><code>9b5dcaf</code></a> Guard read-tree index output paths</li> <li><a href="https://github.com/gitpython-developers/GitPython/commit/96a888f4d782cb2f80452148e48e60ce4af6d541"><code>96a888f</code></a> Check joined short-option values before Git execution</li> <li>Additional commits viewable in <a href="https://github.com/gitpython-developers/GitPython/compare/3.1.57...3.1.58">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>
…ck_dim (#9791) ### Summary of proposed changes Fit the **bsrmv/bsrxmv** general-path wave32 launch geometry to `block_dim`. On wave32 these routes always go through `bsrxmvn_general`, whose `block_dim <= 8` bucket launched `BLOCKSIZE=64, WFSIZE=8` regardless — so for `block_dim` 2 and 4 only `(block_dim/8)^2` of the launched threads did useful work. - Add `block_dim <= 2 → <BLOCKSIZE=4, WFSIZE=2>` and `block_dim <= 4 → <16, 4>` buckets. `BLOCKSIZE == WFSIZE^2` is preserved, so the per-row stride and computed result are unchanged. - These tight-fit buckets are gated on wave32; wave64 parts (with their own specialized small-block paths) fall through to the original `block_dim <= 8 → WFSIZE=8` bucket unchanged. - `WFSIZE` is a logical row-owner group (not the hardware wave), so results are identical on every architecture. **Performance** (gfx1201 / RX 9070): **+80% to +150%** on `block_dim` 2/4; `>= 8` buckets (already bandwidth-bound) untouched. **Correctness**: numerics unchanged; `3811/3811` bsr(x)mv tests pass. **Portability / risk**: wave32-gated; wave64 unchanged. JIRA ID : AISPARSE-621 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…atch count >65536 (#9975) ## Motivation Currently we use the y grid dimension to perform batch computation with SpMM COO segmented algorithm. This means that if the number of batches exceeds 65535 we will get launch errors. This PR fixes that by looping over the batch count in blocks of 32768 batches. JIRA ID: AISPARSE-606
## Motivation It has been decided from management to postpone the inclusion of CCCL 3.0 support. ## Technical Details This PR reverts all commits related to CCCL 3.0. See the list of PRs in the commit list. ## Test Plan Built benchmarks and tests locally to verify nothing was broken. Running CI to get the bigger picture. ## Test Result JIRA ID : ROCM-29174 EXSWSTRHPC-300 --------- Co-authored-by: spolifroni-amd <sandra.polifroni@amd.com>
… awareness (#10409) ## Motivation JIRA ID: AIHPBLAS-1405 InsertCoexecHazardPass wrongly reduced co-exec V_NOP counts when DISABLE_XDL_ARB_STALL was set, but that bit is SQ issue-pacing only that it does not disable SP co-execution. Result: dependent WMMA→VALU pairs got 0 V_NOPs and read stale Matrix D → silent data corruption. Separately, the DAG scheduler didn't model the co-exec gap, so it left gaps for the V_NOP pass to pad instead of filling them with useful VALU. <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details - Fix (counts): counts are now always co-exec-ON (WMMA→WMMA = slots+1, WMMA→VALU = slots); - DAG awareness: the scheduler spaces a dependent b2b WMMA / WMMA→VALU by interleaving independent VALU (only VALU fills a co-exec slot; vmem/SALU keep filling the window for throughput but don't count). When no independent VALU is left, it emits the shortfall as V_NOPs itself. - V_NOP ownership: StinkyRemoveNopPass gains a vNopOnly option and strips stray v_nops before scheduling; InsertCoexecHazardPass now counts existing v_nops and tops up the shortfall instead of strip-and-re-emit, so the scheduler's placed fillers survive. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Result - [x] ctest - [x] tox on machine (exclude existing failed cases) <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Summary Add the `rocke: platform` label to PRs that modify files under: `dnn-providers/hip-kernel-provider/rocke/platform/**` ## Why The rocke project is separated into platform and library, We want to add this platform label to differentiated the two. Project labels are [Strata](http://perf.adc.amd.com:3010/explorer/pull-requests)s primary source of truth, so these will also be used to seperate rocke work in Stata. ## Testing - Confirmed `.github/labeler.yml` parses successfully. - Confirmed the label maps to the intended Rocke platform path. JIRA ID : AICK-1871 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: andrew clark <anclark@amd.com>
## Motivation Widen backward and forward compatibility for FFT client programs. JIRA ID: AIFFT-257 ## Technical Details rocFFT will runtime-compile any kernels it needs on an architecture it wasn't built to know about, but the FFT client programs need kernels built at project build time. Add amdgcnspirv architecture to the client programs to ensure that they will run on any arch supported by the runtime, rather than only those that were specified at build time. Reapplies 0140ec0 which was reverted by 7ac51a7. ## Test Plan Manually built TheRock with one architecture and ran that build's client programs on another architecture. Client programs ran successfully, even with kpack enabled. Manually checked ASAN build of rocFFT ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
- Remove hardcoded prebuilt_stages and baseline_run_id - Add stage_reuse_mode: reuse-stage to enable automatic baseline selection - TheRock now finds commit-compatible baselines and determines which stages to rebuild based on changed files in rocm-libraries ISSUE ID: ROCm/TheRock#3343 --------- Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-4258 Follow-up PR of #10355 ## Motivation gfx1250 CGEMM already supports WMMA+TDM codegen (#10355), and already add some CGEMM+TDM solutions. This PR adds more the CGEMM TDM solutions to logic yamls to improve coverage. <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details Adds WMMA+TDM solutions to NN/NT/TT of the gfx1250 CGEMM logic yamls (`gfx1250_Cijk_*_C_B_UserArgs.yaml`), merged into the existing logic yamls. After this PR each transpose has >=3 TDM solutions (NT=3, NN=4, TT=3, TN=3). p.s. TN already had 3 TDM solutions from #10355 and is left unchanged in this PR. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan Clean build for gfx1250, run `hipblaslt-test` including CGEMM (`--gtest_filter=*matmul*complex*`). <!-- Explain any relevant testing done to verify this PR. --> ## Test Result All passed <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: mengzcai <mengzcai@amd.com>
…quired (#10494) ## Motivation JIRA ID: ROCM-28393 Fix a possible crash (access violation) in hipfftw functional validation tests, due to test-side bugs introduced in #7189. A crash occurs non-deterministically when consecutive tests share the same FFT configuration (allowing cache reuse), and manifests as a NULL-plan dereference inside by `fftw_execute_dft`. Copy operations from cached reference-results to test buffers were also considering an incorrect (possible too-large) batch size. ## Technical Details `hipfftw_functional_validation::SetUp` unconditionally called `initialize_input` and `launch_async_compute` on the `reference_fft_data_t` object, even when its constructor successfully picked up cached results from a previous test. Since `execute_cpu_fft` destroys the FFTW plan after use, re-launching the async compute on cached data passes a NULL plan to `fftw_execute_dft`, causing the access violation. The fix guards both calls behind `needs_input_initialization()` and `needs_computing()`, skipping re-computation when the cache already holds valid results for the test instance of interest. ## Test Plan Current test base suffice. Manually tested on configurations found to fail/crash. ## Test Result Tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
… builds (#10354) Follow-up to #8834. `TestRaggedTensor.LargeOffsetExceedsInt32Max` passed in release but failed in debug builds: it called `getIndex(1)` on a single-batch (B=1) tensor, where `ragged_offset[1] == 2^31` is the one-past-the-end sentinel, not an addressable index. The debug-only bounds check (`#ifndef NDEBUG`) in `Tensor::throwIfOutOfBounds` correctly rejected it and threw; release compiled the check out, so the same call silently passed. The test now uses a two-batch (B=2) tensor so `ragged_offset[1] == 2^31` is an interior, in-bounds base, letting `getIndex(1)` legally return 2^31 and exercise the full addressing path in both build configurations.
## Motivation Optimize Conv for pointwise cases. ## Technical Details Simplify tensor descriptors ## Test Plan conv smoke test ## Test Result passed ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
## Motivation Enable possiblity to compare rocke conv kernels with CK kernels. ## Technical Details - add possiblity to run ck conv kernels - new bench shapes ## Test Plan N/A ## Test Result N/A ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Changes are behind the feature flag HIPKERNELPROVIDER_ENABLE_ROCKE which is defaulted to OFF. JIRA ID: AICK-1742 --------- Co-authored-by: Your Name <you@example.com>
#10264) JIRA ID : ROCM-20223 ## Summary Refreshes RPP documentation to match its current status as a monorepo project delivered via TheRock. - Document monorepo move (ROCm/rpp → ROCm/rocm-libraries) and TheRock delivery for ROCm 7.15+; drop OpenCL references (removed in 3.0.0) - Rebuild the CPU/GPU support matrix from the current API headers; replace the un-diffable PNG with Markdown tables in the README and sync the docs `.rst` table (101 entries, verified in sync) - Bump version to 3.2.0 and consolidate the CHANGELOG unreleased section; wire Doxygen version from CMakeLists - Fix stale links, typos, and remove the outdated `RPP-functions.csv` ## Test plan - [ ] Docs build (Sphinx + Doxygen) renders without warnings - [ ] README Markdown tables and collapsible sections render correctly on GitHub
…0522) I've been monitoring the rocjitsu race-detection sidecar that landed in [PR #9969](#9969). The first sidecars to run after the merge were triggered on [PR #9631](#9631), and both failed before reaching either detector workload: ```text OpenMP runtime not found: .../lib/llvm/lib/x86_64-unknown-linux-gnu/libomp.so ``` This looked familiar because a missing `libomp.so` was one of the reasons the original sidecar was temporarily reverted in [PR #9938](#9938). The original failure and the current failure have the same underlying dependency, but the compiler artifact layout changed: - PR #9969's green CI proof used a compiler artifact with `libomp.so` under: ```text lib/llvm/lib/x86_64-unknown-linux-gnu/libomp.so ``` - Current `develop` uses a newer TheRock revision that includes [TheRock PR #6802](ROCm/TheRock#6802), which reverted host per-target runtime directories. Its compiler artifact places the runtime at: ```text lib/llvm/lib/libomp.so ``` [TheRock PR #7082](ROCm/TheRock#7082) is now re-enabling `LLVM_ENABLE_PER_TARGET_RUNTIME_DIR`. Changing the hard-coded path to the current flat location would therefore regress when the nested layout returns. This PR instead supports both expected layouts explicitly. The sidecar script added the historical host-triple directory to `LD_LIBRARY_PATH`, but it also required `libomp.so` to exist at that exact path. With the current artifact, the parent LLVM library directory is already in `LD_LIBRARY_PATH`, but the hard-coded preflight check exits before the workload can run. This PR replaces that exact-path assumption with a deterministic ordered lookup: 1. `lib/llvm/lib/libomp.so` 2. `lib/llvm/lib/<machine>-unknown-linux-gnu/libomp.so` If neither expected path contains `libomp.so`, the script still fails early and prints both attempted locations. ## Validation Focused validation invoked the real script against temporary ROCm roots while stubbing the later compiler invocation. The observed results were: ```text flat: LIBOMP_PATH=.../lib/llvm/lib/libomp.so lib/llvm/lib appears once in LD_LIBRARY_PATH nested: LIBOMP_PATH=.../lib/llvm/lib/x86_64-unknown-linux-gnu/libomp.so both the root and host-triple directories appear in LD_LIBRARY_PATH missing: OpenMP runtime not found at either expected path: .../lib/llvm/lib/libomp.so .../lib/llvm/lib/x86_64-unknown-linux-gnu/libomp.so layout validation passed: flat, nested, missing ``` ## Issue Tracking Related: AIHPBLAS-3605
…tion (#9648) ## Summary Add the generic `hipsparseSpGEAM` API for sparse matrix-matrix addition, `C = alpha * op(A) + beta * op(B)`, matching the cuSPARSE 13.3 generic `SpGEAM` API. This wraps the staged rocSPARSE `rocsparse_spgeam` on the AMD/ROCm backend and passes through to `cusparseSpGEAM` on the NVIDIA/CUDA backend. - New entry points `hipsparseSpGEAM_createDescr`, `hipsparseSpGEAM_destroyDescr`, `hipsparseSpGEAM_bufferSize`, `hipsparseSpGEAM_nnz`, and `hipsparseSpGEAM`, plus the `hipsparseSpGEAMDescr_t` type and the `hipsparseSpGEAMAlg_t` enum (`HIPSPARSE_SPGEAM_ALG1`). All guarded by `CUDART_VERSION >= 13030` so headers still compile on CUDA toolkits without SpGEAM. - Supports CSR matrices, compute types `HIP_R_32F`/`HIP_R_64F`/`HIP_C_32F`/`HIP_C_64F`, and 32-/64-bit index types. Output C column indices are returned sorted. - API reference documentation (`generic.rst`, `types.rst`, `api.rst`) and the CHANGELOG are updated. A documentation example is provided. ## Testing - Unit tests under `clients/tests` (`spgeam_csr`, `spgeam_csr_bad_arg`) covering data types, index types, and edge cases. - A and B use independent sparsity patterns so the pattern-merge path and the sorted-output-column guarantee are validated against a host `csrgeam` reference. - `alpha`/`beta` zero edge cases are covered. Note: the sparsity pattern of C is the union of A and B and is independent of the scalar values on both backends (this matches cuSPARSE's documented behavior, where zero scalars are not a special case). - Unsupported cases return a well-defined error: transpose operations and non-CSR formats now return `HIPSPARSE_STATUS_NOT_SUPPORTED` in both backends (previously a transpose op was silently accepted on the ROCm backend). - All 96 spgeam tests pass on gfx1201 (ROCm backend). The NVIDIA-side guard mirrors the ROCm logic and uses only public API; it should be confirmed on a CUDA build. JIRA ID : AISPARSE-586 Made with Cursor Made with [Cursor](https://cursor.com) --------- Co-authored-by: Nico Trost <nico.trost@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>
…Bwd/Wrw CK solvers (#10240) JIRA ID: ALMIOPEN-2359 ## Motivation `miopenFindConvolutionBackwardDataAlgorithm()` (and the analogous BackwardWeights find) could hang indefinitely on MI300X whenever three conditions coincided: TF32 data type, the `MIOPEN_CONVOLUTION_ATTRIB_DETERMINISTIC` attribute set on the convolution, and an exhaustive find with an empty perf-db cache. Root cause: an inverted return-value polarity bug in the `is_deterministic` branch of `SetNextValue()` for `PerformanceConfigHipImplicitGemmGroupBwdXdlops` and `PerformanceConfigHipImplicitGemmGroupWrwXdlops`. `NextLinear()` returns `true` on wraparound (search exhausted) and `false` on a normal advance. The old code did `if(!NextLinear(...)) return false;`, i.e. it treated a *normal advance* (`!NextLinear` == `true`) as "exhausted" and a *wraparound* (`NextLinear` == `true`, so `!NextLinear` == `false`) as "keep going" — exactly backwards. The practical fallout: - With exactly one valid CK kernel candidate, the index space never wraps, so the branch never returned `false` and the enumeration spun forever (the observed hang). - With more than one valid kernel, the very first call incorrectly reported "exhausted" and returned `false`, so the search silently visited zero performance configs (under-tuning: the deterministic path never actually explored the kernel list). The unaffected sibling solvers, `PerformanceConfigHipImplicitGemmGroupFwdXdlops` and `PerformanceConfigHipImplicitGemm3DGroupBwdXdlops`, already implement the equivalent deterministic-mode advance correctly using a direct boundary check instead of `NextLinear()`'s return value, which is why only the grouped Bwd/Wrw Xdlops solvers were affected. ## Technical Details - `projects/miopen/src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp` and `.../conv_hip_implicit_gemm_grouped_wrw_xdlops.cpp`: replaced the `if(!NextLinear(0, valid_kernels.size() - 1, index))` idiom in the `is_deterministic` branch of `SetNextValue()` with the same explicit boundary-check idiom already used correctly by the Fwd and 3D Bwd siblings: ```cpp if((index + 1) < static_cast<int>(valid_kernels.size())) { ++index; split_k = 1; kernel_id = valid_kernels[index] + "+1"; return true; } return false; // All kernels exhausted ``` This makes the deterministic branch self-contained (no dependence on `NextLinear`'s wraparound semantics) and consistent with the rest of the codebase. The non-deterministic branch (which still uses `NextTwoPower`/`NextCKSplitkValue`) is untouched. - `projects/miopen/src/include/miopen/conv/solvers.hpp`: added `MIOPEN_INTERNALS_EXPORT` to both affected `SetNextValue()` declarations, matching the annotation already present on the 3D Bwd sibling's `SetNextValue()`. This is what lets the new gtest binary call `SetNextValue()` directly from a host-only test without linking against the CK shared library or requiring a GPU. - `projects/miopen/test/gtest/perf_config_group_deterministic_set_next_value.cpp` (new file): host-only regression test with 6 cases covering, for both Bwd and Wrw: - the hang regression (one valid kernel — must return `false` immediately, not loop), - the under-tuning regression (three valid kernels — must visit index 1 then index 2 before exhausting), - a non-deterministic-branch smoke guard (asserts the untouched branch still visits more than one index/split_k combination and terminates within a bounded iteration count). The tests populate `valid_kernels` directly on the performance config before calling `SetNextValue()`, bypassing the lazy CK-library-backed candidate lookup, so they run entirely on the host with no GPU. ## Risk Assessment **Risk Level: 🟡 Medium** The code change itself is minimal (two near-identical 5-line branch rewrites plus two export annotations), and the affected code path is gated behind the opt-in `MIOPEN_CONVOLUTION_ATTRIB_DETERMINISTIC` attribute, so most callers are unaffected. It is rated Medium rather than Low because it touches core convolution solver performance-config search logic (`conv_hip_implicit_gemm_grouped_{bwd,wrw}_xdlops.cpp`) — exactly the kind of file whose subtle polarity/boundary bugs are hard to catch by inspection (as this bug itself demonstrates) and which is exercised by every grouped Bwd/Wrw Xdlops find/tune call. ### Impacted Components - `miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupBwdXdlops::SetNextValue` (grouped backward-data CK Xdlops solver) - `miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupWrwXdlops::SetNextValue` (grouped backward-weights CK Xdlops solver) - `projects/miopen/src/include/miopen/conv/solvers.hpp` (export visibility only, no behavioral change) - New test target: `perf_config_group_deterministic_set_next_value.cpp` ### Potential Side Effects - Deterministic-mode `Find`/tuning calls for grouped Bwd/Wrw Xdlops convolutions will now actually enumerate all valid kernel candidates instead of hanging (single candidate) or silently enumerating zero candidates (multiple candidates). Any downstream code or test that was inadvertently relying on the previous zero-candidates-enumerated behavior (e.g. assuming the deterministic path always falls back to a default/heuristic config) could see a different kernel selected. - No change to the non-deterministic branch, to `IsValidValue()`/`IsValid()`, or to any other solver. - The added `MIOPEN_INTERNALS_EXPORT` annotations only affect symbol visibility for internal builds/tests; they do not change runtime behavior. ### Mitigation Steps - The fix reuses an idiom already proven correct and shipping in two sibling solvers (`GroupFwdXdlops`, `3DGroupBwdXdlops`) rather than inventing new logic, minimizing the chance of a new class of bug. - Added a dedicated, fast, host-only regression test suite that directly exercises the exact hang and under-tuning scenarios, so any future regression on this branch is caught in CI without needing GPU hardware or a TF32/MI300X repro environment. - Revert plan: the change is a self-contained 2-file, ~10-line functional diff plus a new test file; reverting `conv_hip_implicit_gemm_grouped_{bwd,wrw}_xdlops.cpp` to the prior `NextLinear`-based branch (and dropping the new test) fully restores prior behavior if needed. ## Test Plan Added `projects/miopen/test/gtest/perf_config_group_deterministic_set_next_value.cpp`, a new host-only gtest file with 6 test cases (no GPU or CK shared library required): - `CPU_PerfConfigGroupBwdDeterministic_NONE.HangRegressionSingleKernel` - `CPU_PerfConfigGroupBwdDeterministic_NONE.UnderTuningThreeKernels` - `CPU_PerfConfigGroupWrwDeterministic_NONE.HangRegressionSingleKernel` - `CPU_PerfConfigGroupWrwDeterministic_NONE.UnderTuningThreeKernels` - `CPU_PerfConfigGroupBwdNonDeterministic_NONE.VisitsMultipleIndicesAndSplitKs` - `CPU_PerfConfigGroupWrwNonDeterministic_NONE.VisitsMultipleIndicesAndSplitKs` The hang-regression cases use a single valid kernel and assert `SetNextValue()` returns `false` on the first call (with a bounded iteration cap of 10 as a safety net against a true infinite loop during test execution). The under-tuning cases use three valid kernels and assert the config advances through index 1 and index 2 (with correct `split_k`/`kernel_id`) before exhausting on the fourth call. The non-deterministic smoke test guards the untouched branch, asserting it still visits more than one index and more than one split_k value within a 500-iteration cap. ## Test Result All 6 new tests were built and run locally and pass. The fix was additionally verified to be meaningful (not a no-op) by temporarily reverting the `SetNextValue()` change and re-running the suite: with the reverted (buggy) code, the corresponding tests fail fast via the bounded-iteration-cap assertion (`ASSERT_LT(iterations, 10)` / `ASSERT_LT(iterations, 500)`) rather than hanging the test process, confirming the test suite correctly detects both the hang and under-tuning regressions. No GPU hardware or TF32/MI300X environment was needed to build or run these tests, since they call `SetNextValue()` directly on a host-constructed `ProblemDescription` with a manually-populated `valid_kernels` list. Co-authored-by: Claude <noreply@anthropic.com>
… test artifact (#10517) The `HIPTHREADS_COPY_TO_BUILD` step blanket-copies the source tree into the packaged test artifact, including `examples/sparse-mat-mul/data/matrices.tgz`. The Windows builder has git-LFS, so it ships the real ~930MB file; the Linux container doesn't, so it ships a pointer, hence the 893MB vs 272KB size gap. Exclude `matrices.tgz` and `shallow_water1.mtx` from the copy so only the committed non-LFS `test_general.mtx` (the file the tests actually use) ships. This affects the artifact only; `.gitattributes` still keeps both as LFS pointers in the repo for users who `git lfs pull`. Closes ROCm/TheRock#7171
) JIRA ID : N/A Close #10467 ## Motivation `ENABLE_CLANG_TIDY` defaults to OFF, but the `CMAKE_MODULE_PATH` append, `include(ClangTidy)` and `add_clang_tidy_custom_target()` all ran unconditionally. `ClangTidy.cmake` lives in the superbuild's top-level `cmake/modules`, two levels above the subproject, so configuring stinkytofu depended on a path outside its own tree even with the feature disabled. Any checkout that does not materialise the top-level `cmake/` directory fails to configure: ``` CMake Error at CMakeLists.txt:366 (include): include could not find requested file: ClangTidy CMake Error at CMakeLists.txt:367 (add_clang_tidy_custom_target): Unknown CMake command "add_clang_tidy_custom_target". ``` The second error is a cascade of the first, so the reported failure is an unknown command rather than a missing file — disproportionately confusing to diagnose. ## Technical Details Moved all three statements inside `if(ENABLE_CLANG_TIDY)`. `include(ClangTidy OPTIONAL)` was deliberately not used: it would silently skip the include and defer the failure to `add_clang_tidy_custom_target()` with a worse message. When clang-tidy is explicitly requested, a missing module should stay a hard error. ## Test Plan Standalone configure of `shared/stinkytofu` with the guard applied, both with the option off (default) and on, checking that clang-tidy targets appear only when requested. ## Test Result ``` cmake -S shared/stinkytofu -B /tmp/st-off -G Ninja -DROCM_PATH=$ROCM_PATH exit=0 -- Build files have been written to: /tmp/st-off cmake -S shared/stinkytofu -B /tmp/st-on -G Ninja -DROCM_PATH=$ROCM_PATH -DENABLE_CLANG_TIDY=ON exit=0 -- Build files have been written to: /tmp/st-on ``` `ninja -t targets | grep -c tidy` → **0** with the option off, **3** with it on. The original failure was reproduced before the fix on a checkout whose `cmake/` directory was not materialised. Not run: full stinkytofu build/test suite (configure-only change), Windows. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level None. Configure-time only, and strictly narrows when the clang-tidy block executes. With `ENABLE_CLANG_TIDY=ON` behaviour is unchanged; with it off the block simply no longer runs.
github-actions
Bot
added
project: rocsparse
project: rocthrust
github actions
project: hipdnn
project: miopen-provider
project: hipblaslt-provider
shared: primbench
ci:hipsparselt-fast
project: hip-kernel-provider
project: hipdnn-integration-tests
shared: stinkytofu
project: rpp
project: rocalution
rocke: platform
## Summary Declares `HIPDNN_ENABLE_KERNEL_INGESTOR` (default `OFF`) in both `projects/hipdnn/CMakeLists.txt` and `dnn-providers/hip-kernel-provider/CMakeLists.txt`, matching the variable TheRock's `FLAGS.cmake` forwards to both subprojects ([ROCm/TheRock#7194](ROCm/TheRock#7194)). This reserves a single option ahead of the parallel follow-on stories (ALMIOPEN-2400, ALMIOPEN-2401, ALMIOPEN-2402) that will land the actual `#ifdef`/`if()`-gated kernel-ingestor code, so those PRs don't each independently declare — and risk conflicting on — the same option. See [RFC 0017: Universal Kernel Descriptors](projects/hipdnn/docs/rfcs/0017_UniversalKernelDescriptor.md) for the platform this gates. JIRA ID : ALMIOPEN-2399 ## Risk Assessment Minimal risk. Adds an inert, default-OFF `option()` in two CMakeLists files with no gated code yet; no existing build path, target, or test is affected. ## ASIC Coverage No ASIC impact. Declaration-only change with no compiled/runtime behavior difference at any flag value. ## Testing Summary - Standalone CMake configure of `projects/hipdnn` and `dnn-providers/hip-kernel-provider` with `-DHIPDNN_ENABLE_KERNEL_INGESTOR=ON`, confirming the option registers in `CMakeCache.txt` as `ON` before either configure hits its (pre-existing, environment-only) missing-dependency errors unrelated to this change. ## Testing Checklist - [x] CMake configure (hipDNN) - `cmake -S projects/hipdnn -B build -DHIPDNN_ENABLE_KERNEL_INGESTOR=ON` - Status: Passed (option registers `ON` in cache) - [x] CMake configure (hip-kernel-provider) - `cmake -S dnn-providers/hip-kernel-provider -B build -DHIPDNN_ENABLE_KERNEL_INGESTOR=ON` - Status: Passed (option registers `ON` in cache) - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - Adds `option(HIPDNN_ENABLE_KERNEL_INGESTOR "Enable the generic kernel ingestor" OFF)` to `projects/hipdnn/CMakeLists.txt`, after the existing option block. - Adds the same option to `dnn-providers/hip-kernel-provider/CMakeLists.txt`, alongside `HIPKERNELPROVIDER_ENABLE_ROCKE`.
All workflows that call setup_multi_arch.yml need actions:read permission for the stage reuse feature to query GitHub API for baseline workflow runs. ISSUE ID: ROCm/TheRock#3343 Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
The two logic files added in c4c1368 were written with the trailing list-format elements (RangeLogic, TileSelectionIndices, PerfMetric and the library-type attribute) nested inside ExactLogic instead of being emitted as top-level keys. TensileCreateLibrary failed with KeyError: 'LibraryType' in prepareLibraryLogicDict, breaking the gfx1250 device-library build. Dedent the ExactLogic table by one sequence level and emit RangeLogic, TileSelectionIndices, PerfMetric and LibraryType as top-level keys, so both files match the shape of every other dict-format logic file. Tuning data is unchanged: 47 solutions / 49 exact entries (NT) and 2 solutions / 2 exact entries (NN), indices contiguous with no dangling references. All 209 gfx1250 logic files now parse. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Equality tuning for 4K x 4K x 65536 shape. Added 4 liblogic yaml files for each BBS_TN, F8F8S_TN, MX_F8BS_TN, MX_F4BS_TN. Tested with tensilelite and hipblaslt-bench verification run. Merged with --admin for the BKC release code cut-off. Math CI Summary and pre-commit green; the reds in TheRock CI Summary are #10522 (rocjitsu libomp.so, fixed on develop after this run was created) and ROCm/TheRock#7170 (30-min amd-llvm clone timeout on gfx125X, with a same-run green gfx950 build as control). Rationale in the PR thread. JIRA ID : AIHPBLAS-4151
An error occurred while trying to automatically change base from
users/minsukim/gfx1250_equality_maf_20260731
to
develop
August 10, 2026 23:26
…10557) ## Motivation JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-4230 LocalRead could carry multiple LDS memory tokens (for example, TDMSplit both-halves reads), but its comment only printed a single token (sync LDSx). This made generated asm/debug output misleading and harder to correlate with actual dependency metadata. The change makes comments list all attached tokens, improving traceability and reducing confusion when validating barrier/wait behavior. ## Technical Details <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan CI ## Test Result ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
…1250_equality_bbstn_20260810 # Conflicts: # projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/Equality/gfx1250_Cijk_Alik_Bljk_BBS_BH_UserArgs.yaml
Contributor
|
This seems to touch a lot of infrastructure issues; can you clean this up please? It doesn't seem to be just touching hipblaslt. |
Contributor
|
This pull request has been inactive for 25 days and will be marked as stale. If you would like to keep this PR open, please:
This PR will be automatically closed in 5 days if no further activity occurs. |
Contributor
|
This pull request has been automatically closed due to inactivity (30 days with no updates). If you'd like to continue working on this, feel free to reopen the PR or create a new one. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds gfx1250 Equality logic tuned for BF16 TN (
Cijk_Alik_Bljk_BBS_BH).Stacked on top of #10250 (
users/minsukim/gfx1250_equality_maf_20260731), which introduces the gfx1250Equality/logic directory. This PR is based on that branch so the diff shows only the BBS TN delta; it will retarget todeveloponce #10250 merges.Status
🚧 Draft — placeholder branch. Currently contains an empty commit only. The tuned
gfx1250_Cijk_Alik_Bljk_BBS_BH_UserArgs.yamlwill be pushed here.Files (planned)
projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/Equality/gfx1250_Cijk_Alik_Bljk_BBS_BH_UserArgs.yamlTest plan
🤖 Generated with Claude Code