[hipSPARSELt] monorepo test by jayhawk-commits · Pull Request #5 · ROCm/rocm-libraries

jayhawk-commits · 2025-04-30T02:31:15Z

No description provided.

This reverts commit 82a24bb.

…ion (#5) * added addition unit tests and implement missing function in match * updated changelog

…edge dispatch (#7064) ## Motivation Fix correctness failures in the `UseSubtileImpl` NonEdge store path for gfx950 BF16 and MXFP4 subtile kernels. These failures were caused by several interrelated bugs in the interleaved (GLS=0) store codegen that manifested when M-guard branches skipped stores at runtime, and by an overly strict edge/NonEdge dispatch check that forced subtile-aligned workgroups into the unoptimized edge path. ## Technical Details **1. vmcnt hazard in interleaved B1 NonEdge store path** (`GlobalWriteBatch.py`) On gfx950 (`SeparateVscnt=False`), loads and stores share the same vmcnt counter. The interleaved store path computed `vmcnt = vlcnt + vscnt` where `vscnt = self.storesIssued`. When M-guard branches skip stores at runtime, the actual vmcnt counter has fewer outstanding operations than codegen assumed, making `s_waitcnt vmcnt(N)` too permissive — the hardware doesn't wait long enough for C-loads to complete before fmacs consume them. Fixed by setting `vscnt = 0` for `UseSubtileImpl` with `GroupLoadStore=False`. This is more conservative but correct regardless of how many stores were skipped at runtime. **2. SrdD increment skipped by M-guard branch** (`GlobalWriteBatch.py`) The SrdD `incToNextRow` was emitted inside the M-guard-skippable region of the store loop. When the M-guard branch skipped the last store in an N-group, the SrdD increment was also skipped, causing subsequent N-group stores to write to the wrong row address. Fixed by deferring the increment and emitting it after the N-group end label. **3. N-group end label placed after next N-group's fmacs** (`GlobalWriteBatch.py`) When `GroupLoadStore=False`, fmacs and stores are interleaved in the same module. The fmacs for N-group K+1 were emitted before the N-group K end label was placed. The M-guard branch (targeting the end label) would skip both the last store of N-group K *and* the fmacs for N-group K+1, leaving the K+1 accumulators at zero. Fixed by flushing the pending N-group end label and deferred SrdD increment before emitting the next N-group's fmacs. **4. Paired store blockIdxM guard mismatch** (`GlobalWriteBatch.py`) `SubtileMGuard` counts valid M-blocks in `MatrixInstM` (16-row) units, but `blockIdxM` for paired and orphan stores was computed in 32-row units (`(tt0-1)//2` or `(tt0*16)//32`). This mismatch caused incorrect OOB guard decisions. Fixed by using `blockIdxM = tt0` (16-row index) to match MGuard units. **5. Scalar fallback for partial paired stores** (`GlobalWriteBatch.py`) Added a scalar `dwordx2` fallback when only the lower M-block in a pair is valid (`MGuard > tt0-1` but not `MGuard > tt0`). Previously, the paired `dwordx4` store would execute for both blocks even when the upper block was OOB. This case arises when the tile remainder has an odd number of `MatrixInstM`-sized blocks (e.g., remainder=48 = 3 blocks of 16 rows). **6. sba=1 orphan store for large macro tiles** (`GlobalWriteBatch.py`) When `MIWaveTile[0]` is large enough that batch boundaries split an (sba=0, sba=1) pair, the sba=1 element had no partner and was silently dropped. Added scalar store handling for this orphan case. **7. Relaxed edge/NonEdge dispatch alignment** (`KernelWriterAssembly.py`) The `checkIsEdgeSubtile` M-dimension alignment was `waveGroupM` (e.g., 48 for MIWT3), requiring the tile remainder to be a multiple of the full wave group height. Reduced to `MatrixInstM` (16), so any remainder that is a multiple of 16 rows takes the optimized NonEdge path. The NonEdge path's MGuard + scalar fallback (fix #5) handles partial wave groups correctly. All changes are guarded by `UseSubtileImpl` — no impact on non-subtile kernels. ## Test Plan - Run `subtile_bf16.yaml` (BF16 BBS/BSS, multiple MIWT configs, SK3, PGR0/PGR2) with tile-aligned sizes (M % 32 == 0, N % 16 == 0) - Run `subtile_mxfp4.yaml` (MXFP4 F4BS/F4HS/F4SS, bias, activations, ScaleAlphaVec, PGR0/PGR2) with tile-aligned sizes (M % 32 == 0, N % 32 == 0) - Verified edge stores are not exercised for tile-aligned sizes by temporarily disabling edge path stores and confirming all tests still pass - Verified with rocgdb breakpoints that MGuard/NGuard values, C-load data, and accumulator values are correct at N-group boundaries ## Test Result - `subtile_bf16.yaml` (tile-aligned sizes): all tests PASSED - `subtile_mxfp4.yaml` (tile-aligned sizes): all tests PASSED - Edge-stores-disabled verification: all tests PASSED (confirms tile-aligned sizes use NonEdge path exclusively) ## 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 <noreply@anthropic.com>

## Motivation https://amd-hub.atlassian.net/browse/AIHPBLAS-1467 ## Technical Details Fixed multiple issues preventing TensileLight from correctly generating and executing kernels when UseBeta=false (beta parameter not used in GEMM operations). Enabled bounds checking validation to work correctly with this configuration. Files Modified 1. Tensile/KernelWriterAssembly.py Issue: KeyError when accessing Beta SGPR register when UseBeta=false Fix: Added conditional check before accessing Beta SGPR if kernel["ProblemType"]["UseBeta"]: moduleExternalArgs.addComment("Read Beta") moduleExternalArgs.addModuleAsFlatItems(self.externalArgLoader.loadAllKernArg( self.sgprs["Beta"], "KernArgAddress", self.states.numSgprBeta)) 2. Tensile/SolutionStructs/Problem.py Issue: UseBeta serialized as integer (0/1) instead of boolean in YAML, causing C++ parser errors Fix: Ensure UseBeta is always stored as boolean self.state["UseBeta"] = bool(self.state["UseBeta"]) 3. client/src/ReferenceValidator.cpp Issue #1: Buffer allocation check didn't verify if buffer pointer was valid Fix: Check both size and pointer validity // Only skip reallocation if size matches AND buffer is valid if(m_cpuResultBufferSize == bytes && m_cpuResultBuffer.get() != nullptr) return; Issue #2: hipFree compiler warning about nodiscard attribute Fix: Cast return value to void in lambda deleter uint8_t* buffer; HIP_CHECK_EXC(hipHostMalloc((void**)&buffer, bytes, 0)); m_cpuResultBuffer.reset(buffer, [](uint8_t* p) { (void)hipFree(p); }); Issue #3: Attempting to validate null/empty tensors Fix: Skip validation for null pointers or zero-sized tensors // Skip validation if pointers are null or maxElements is 0 if(resPtr == nullptr || refPtr == nullptr || result.maxElements[i] == 0) { if(Debug::Instance().printTensorInfo()) std::cout << "Skipping validation for tensor " << tensor.getName() << std::endl; continue; } Issue #4: Trying to copy padding bytes from output tensors that don't have padding Fix: Only use maxElement for input tensors // For output tensors, don't use maxElement with padding if(boundsCheck == BoundsCheckMode::NaN && !tensor.isOutput()) elementsToCopy = maxElement; Issue #5: Bounds checking validation on output tensors without padding buffers Fix: Skip bounds checking for output tensors // Only check bounds for input tensors (output tensors don't have padding buffers) if(boundsCheck == BoundsCheckMode::NaN && !tensor.isOutput()) 4. client/src/DataInitialization.cpp Issue: hipMemcpy with null pointers causing runtime errors Fix: Added null pointer check void* copyInputBuffers(const TensorDescriptor& descriptor, void* dst, void* src, size_t totalElements, hipMemcpyKind kind) { // Skip copy if no elements to copy or if pointers are null if(totalElements > 0 && dst != nullptr && src != nullptr) { HIP_CHECK_EXC(hipMemcpy(dst, src, descriptor.elementBytes() * totalElements, kind)); } return dst; } 0d6cd23 ## Test Plan NA ## Test Result ``` ========================================================================================== 105 passed, 83 skipped, 1 warning in 1040.87s (0:17:20) ========================================================================================== ``` ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Signed-off-by: pdhirajkumarprasad <dhirajp@amd.com>

…, move inspection tool - Remove reviewer name from RFC body (line 543) - Remove generate_diagrams.py from repo (keep locally) - Remove dead graph_level_correctness diagram code from script - Move Bundle Inspection Tool from Detailed Design to Future Work item #6 (was marked v2/not-v1 but sitting in Detailed Design — confusing) - Add metadata sidecar to Future Work item #5 - Rewrite infrastructure table with consistent Read/Split/Compare pattern Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

…o Future Work - Change "partially working for batchnorm" to "initial infrastructure is in place" - Add Future Work item #5: external data validation (Python-only comparison of client-submitted bundles against golden references) Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

- Generator auto-derives the output path (Operation/Layout/DataType) from graph content — developer supplies only tier and bundle name - Add CLI example: --tier smoke --name Small → full path computed - Update Folder Convention table: Path is auto-derived, not manual - Future Work #5: auto-tier classification based on tensor element counts, matching getSmall/getMedium/getLargeEdge/getLargeStress Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

Refactor pass addressing sebvince review comment #5 ("scale gating should not live in the InstructionEmitter") and partially #6. No behavioral change — R=1 AND R=2 codegen are sha256-bit-identical to pre-refactor. Design change: the *scheduler* now owns the per-iter cadence decision. LogicalScheduler._assign_ui_slots walks the emitted schedule once and tags each scale op with the body-copy slot it should fire in: - scale gr / gr_inc → ui_slot=0 (DTL + SRD advance + LW swap at START of the R-period) - scale lr_inc → ui_slot=R-1 (LDS read-side swap AFTER all R-period scale reads) - data ops / mfma / lr → ui_slot=None (fire on every body copy) InstructionEmitter.populate becomes generic — it has no idea what "scale" means; it just skips ops whose ui_slot does not match ``unroll_iter % R``. Under R==1 every op stays ui_slot=None and populate is bit-identical to the legacy no-gating path. Removed: - InstructionEmitter._scale_op_gated_out (review #5; replaced by generic em.ui_slot check) Kept (with documented rationale): - LogicalScheduler.insert_gr_lr_inc R>1 postOp branch (review #6 asked to delete; this is actually about MT-transition partition- handoff LR placement, not per-ui gating — without it, the test_scheduler_R2_FP8_MT256_partN2_symmetric_mxsa_lr assertions on partition-0 MXSA postOp anchoring fail. The full asymmetric tile-map refactor would make this fall out naturally, but that change is out of scope here.) - InstructionEmitter._gate_last_iter + strip_prefetch (review #7 flagged vmcnt under stripped scale gr; the vmcnt logic counts per-tensor wait_gr_counts set at schedule time and is unaffected by where the per-ui gating decision lives. The OOB SrdMXSA fix on last iter remains a real K-bound concern, independent of the gating-location refactor.) - emit_lr_inc / emit_gr_inc are now truly "fire every body copy" on the data side (no per-ui logic at all in the emitter). Verification: - Unit tests: 111 passed (no change). - R=1 subtile_mxfp8.yaml codegen: sha256-bit-identical (2/2 .s). - R=2 subtile_mxfp8_mt256.yaml codegen: sha256-bit-identical (2/2 .s, all 7 sizes accepted: PGR=0 × {256,512,1024} + PGR=2 × {256,512,1024,4096}). Co-authored-by: Cursor <cursoragent@cursor.com>

Registry_DqDkDv_ReturnsNullForUnregistered probes d96 to assert findVariant() returns null. If a d96 variant is ever registered, that probe silently passes for the wrong reason. Add a loop over ALL_DQDKDV_VARIANTS asserting no d96 dqdkdv entry exists, so the premise fails loudly and forces a new probe instead. Also fix the file-header NOTE: drop the dead dqdkdv_spec.hpp:200-204 line reference and explain that block_n0 is identical (128) for d64 and d128, so hdim_q/hdim_v gate tile selection (verified by the device static_assert). Collapse the four d64 test initializers per clang-format-18. Addresses review findings #5, #6, #9, #10 on PR #7804. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

…E verdict, perf median Address remaining Copilot inline review comments and improve_advice.pdf gaps: cpp_identifier_oracle.cpp: - Add explicit #include <cstdlib> (was relying on transitive includes; Copilot line 29) check_identifier_parity.py: - _ensure_oracle() now compares binary mtime against max(oracle_src, kernel_key.hpp) mtime, not just oracle_src. Prevents stale binary from silently returning wrong identifiers after an in-place kernel_key.hpp edit (Copilot line 103) drive_codegen.py: - After codegen, assert exactly one header was emitted and that the expected identifier appears in its filename. Returns non-zero on either violation so CI catches misnamed or duplicate headers early (improve_advice #2) check_parity.py: - run_te_benchmark() returns an explicit verdict dict {verdict, tflops, rc} instead of Optional[Dict]. Missing CSV → verdict=SKIPPED (TE skipped the size), not a failure (Copilot line 323) - _adjudicate_numerical() treats TE SKIPPED as a per-size skip, not a failure; uses the new verdict field rather than inferring pass/fail from CSV presence - _adjudicate_performance() collects _PERF_RUNS=10 harness invocations per size and compares the median TFLOP/s against the TE baseline; suppresses GPU clock transients and OS scheduler jitter (improve_advice #5) - --perf-tol default changed from 0.10 → 0.02 (2%) per improve_advice #5 spec - --perf-runs flag added to override the 10-run default - import statistics added (used by median computation) harness.cpp: - Verify loop now counts pass/fail elements and prints "N/total (X%)" summary - First 10 mismatches printed to stderr on failure to aid diagnosis without flooding stdout (improve_advice #4) - Added #include <tuple> for structured binding in mismatch vector README.md: - Added Measurement Methodology section documenting warmup=3, repeat=20, GPU timer, 10-run median, 2% tolerance with rationale (improve_advice #6) All 36 unit tests still pass (python3 -m pytest test_te_to_dispatcher.py -v). Co-Authored-By: Claude Sonnet 4 <noreply@anthropic.com>

…patcher) Add Key-invariant #5: depmap + ctest -N + the CI build must come from the same cmake configure, since selection coverage == configure coverage. Names the optional, separately-gated components this affects — rocm_ck (CK_ENABLE_ROCM_CK), codegen/composable_kernel_host (CK_USE_CODEGEN, gfx9; tests codegen_test_*), and dispatcher — and the codegen embed/hiprtc runtime-dependence nuance. codegen/test needs no special tooling (it's a normal C++ component); it's covered when the depmap is generated from the CI configure. Docs only. Generated-by: Claude Code (claude-sonnet-4-6)

…#3, #5, stacked on #1, #2, #5] Generalize the fast path to support double-precision accumulation: - Template ShadowBuffer<AccumT>: storage, pointer, and element access are all AccumT. Float/Double inputs zero-copy when AccumT matches; sub-float types go through float then widen. - Template loadTo<AccumT, SrcType> and storeFrom<AccumT, DstType> (renamed from loadToFloat/storeFromFloat). - Rename solveCPUFastInF32 → solveCPUFast<AccumT, MathOpAccumT>: all tile registers, inner reduction, epilogue, alpha/beta extraction, bias reading, and activation args use AccumT. - Add Double to isFastPathEligible's supported input/output types. - SolveGemmCPU dispatch: route Double to solveCPUFast<double>. - Add 10 f64 fast-path tests (transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector). Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

The fast and slow CPU GEMM paths disagreed about what one-sided MX (mxBlockA>0 with mxBlockB==0, or vice versa) means: the fast path treated the missing side as scale=1.0 while the slow path silently ignored MX entirely. Production code always supplies both scales, so reject one-sided MX in both paths to keep them consistent. isFastPathEligible now rejects when exactly one of mxBlockA/mxBlockB is > 0; runGemm splits the scalar --mxBlock into mxBlockA/mxBlockB locals and errors out symmetrically. mxBlockA != mxBlockB (both > 0) remains allowed and is covered by review #5.

…eview #5) Add per-side MX block flags to the driver to exercise mxBlockA != mxBlockB end-to-end. --mxBlock is kept as a "set both" shortcut; combining it with either per-side flag is rejected. One-sided MX (only A or only B > 0) is also rejected at the driver, matching the reference-path rejection from review #1. The columnMajorGemm golden reference now steps the inner reduction by min(mxBlockA, mxBlockB) so each segment has constant (sa, sb); both blocks are powers of 2 so one divides the other and the step is valid. Asymmetric MX (mxBlockA != mxBlockB) is rejected at the driver for the slow path: the production slow path's MX inner loop applies a single scale per max(mxBlockA, mxBlockB)-sized segment, which collapses the smaller-blocked side's per-segment scales and produces wrong results; fixing the production slow path is out of scope. Tests cover fast path A32/B64 and A64/B32 (K=192 = lcm), plus negative cases for the combinations the driver rejects.

Records the 5 approaches considered for symbolic-vs-numeric register robustness: 1. Name-resolution table (brittle, complex) 2. Symbolic-only normalization (doesn't solve actual problem) 3. Numeric-only resolution (assembly-time dependency) 4. Equivalence-class comparison (loses precision) 5. Render-string identity (matches GPU view, robust) And the rationale for picking #5. Documents the known limitation: same logical reg with different identifiers across captures still differs. Doesn't arise in practice because both captures consume the same writer state; future work if needed would add approach #1 (name-resolution table) on top.

#5) Completes the coverage-gap inventory for ScheduleCapture._reads. Pins the LR LDS-address gap (DSLoad src is LocalReadAddrA, modified by LRS VXorB32 — invisible RAW today) and the LW LDS-address gap (DSStore dstAddr is LocalWriteAddrA, same VXorB32 producer pattern). Both fail loudly when Sub-task 10's wrapper-based extractors close the gap.

Adds _iter_note(producer, consumer) in ScheduleCapture.py: returns " (of next iteration)" when consumer.position.loop_index == producer.position.loop_index + 1. Generalizes the prior MissingWaitFailure inline check (which hardcoded BODY_LABEL_TO_LOOP_INDEX[ML_PREV] -> [ML]) to any i -> i+1 boundary; loop_index is the canonical cross-body iteration counter so the numeric +1 test is the right discriminator. MissingWaitFailure (#2) refactored to use the helper. WaitTooLateFailure (#4), WaitInsufficientFailure (#5), and MissingBarrierFailure (#6) now also append the suffix when the producer/consumer pair crosses an iteration boundary. Suffix attaches right after the consumer's `@ idx=N` mention so the message reads: MFMA[name] @ idx=10 (of next iteration) is guaranteed by an SWaitCnt @ idx=12 ... MFMA[name] @ idx=10 (of next iteration)'s producer LRA0 @ idx=5 ... ... between the SWaitCnt and GRA @ idx=2 (of next iteration). 3 new cross-iter pinning tests + 3 same-iter regression assertions (`assert "(of next iteration)" not in msg`) so a future regression that incorrectly fires the suffix on same-iter pairs is caught.

Adds _node_with_pos(node, capture) — combines _node_label (per-category- stream [N] index, plain MFMA omits) with bare '@ idx=M' position render. Single helper for the canonical Failure node reference shape, replacing three different prior styles: - #1, #2: manual `_node_label + " " + "@ idx=N"` concatenation - #4, #5: `category[name] format_position(...)` (rendered the FULL name inside brackets, e.g. `LRA0[LRA0[0]]`, plus a cross-category list suffix that duplicates [N]'s purpose) - #6: `category format_position(...)` (no brackets) All five formatters now route through _node_with_pos. Plain MFMA stays bracket-less per _node_label's MFMA discriminator; PackMFMAs (categories PackA*/PackB*) keep [N] because CMS reschedules them. Waits in #5 stay as bare '@ idx=N' — the surrounding 'SWaitCnt' word already names the kind; rendering the SYNC category as `SYNC[N] @ idx=M` would just duplicate that. Skipped: - #10 SCCConflict: brackets carry rocisa class name (e.g. [SCSelectB32]) not [N] index; semantic conflict, separate audit. - #7 WrongInterleaving / #8 TimingTooClose: use `name` field for Pack identity (MiddlePack_a/_b/_c, CVT0_a/_b); replacing with [N] would lose the a/b/c discriminator. - #13 ConstraintViolation: slated for deletion in bead `pcz`. 3 new pinning tests verifying [N] actually appears when capture is given (one per Failure: #4 LRA0[1], #5 LRA0[1] + plain MFMA bracket-less, #6 GRA[1] in trailing reference). Test count: 564 passed (+3).

…#3, #5, stacked on #1, #2, #5] Generalize the fast path to support double-precision accumulation: - Template ShadowBuffer<AccumT>: storage, pointer, and element access are all AccumT. Float/Double inputs zero-copy when AccumT matches; sub-float types go through float then widen. - Template loadTo<AccumT, SrcType> and storeFrom<AccumT, DstType> (renamed from loadToFloat/storeFromFloat). - Rename solveCPUFastInF32 → solveCPUFast<AccumT, MathOpAccumT>: all tile registers, inner reduction, epilogue, alpha/beta extraction, bias reading, and activation args use AccumT. - Add Double to isFastPathEligible's supported input/output types. - SolveGemmCPU dispatch: route Double to solveCPUFast<double>. - Add 10 f64 fast-path tests (transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector). Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

…#3, #5, stacked on #1, #2, #5] Generalize the fast path to support double-precision accumulation: - Template ShadowBuffer<AccumT>: storage, pointer, and element access are all AccumT. Float/Double inputs zero-copy when AccumT matches; sub-float types go through float then widen. - Template loadTo<AccumT, SrcType> and storeFrom<AccumT, DstType> (renamed from loadToFloat/storeFromFloat). - Rename solveCPUFastInF32 → solveCPUFast<AccumT, MathOpAccumT>: all tile registers, inner reduction, epilogue, alpha/beta extraction, bias reading, and activation args use AccumT. - Add Double to isFastPathEligible's supported input/output types. - SolveGemmCPU dispatch: route Double to solveCPUFast<double>. - Add 10 f64 fast-path tests (transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector). Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> (cherry picked from commit 1943c7c)

…#3, #5, stacked on #1, #2, #5] Generalize the fast path to support double-precision accumulation: - Template ShadowBuffer<AccumT>: storage, pointer, and element access are all AccumT. Float/Double inputs zero-copy when AccumT matches; sub-float types go through float then widen. - Template loadTo<AccumT, SrcType> and storeFrom<AccumT, DstType> (renamed from loadToFloat/storeFromFloat). - Rename solveCPUFastInF32 → solveCPUFast<AccumT, MathOpAccumT>: all tile registers, inner reduction, epilogue, alpha/beta extraction, bias reading, and activation args use AccumT. - Add Double to isFastPathEligible's supported input/output types. - SolveGemmCPU dispatch: route Double to solveCPUFast<double>. - Add 10 f64 fast-path tests (transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector). Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

…#3, #5, stacked on #1, #2, #5] Generalize the fast path to support double-precision accumulation: - Template ShadowBuffer<AccumT>: storage, pointer, and element access are all AccumT. Float/Double inputs zero-copy when AccumT matches; sub-float types go through float then widen. - Template loadTo<AccumT, SrcType> and storeFrom<AccumT, DstType> (renamed from loadToFloat/storeFromFloat). - Rename solveCPUFastInF32 → solveCPUFast<AccumT, MathOpAccumT>: all tile registers, inner reduction, epilogue, alpha/beta extraction, bias reading, and activation args use AccumT. - Add Double to isFastPathEligible's supported input/output types. - SolveGemmCPU dispatch: route Double to solveCPUFast<double>. - Add 10 f64 fast-path tests (transpose combos, Beta, Bias, AllFeatures, TN_AllFeatures, ScaleAB Scalar/Vector). Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> (cherry picked from commit 1943c7c)

jayhawk-commits and others added 3 commits April 29, 2025 22:18

Add workflow to auto-label incoming PRs

cdf6925

Fixes to auto-labeler

f3a7c88

[hipSPARSELt] monorepo test

82a24bb

jayhawk-commits force-pushed the develop branch from f3a7c88 to 7669408 Compare April 30, 2025 02:40

jayhawk-commits added 4 commits April 29, 2025 22:41

Merge branch 'develop' into joseph-monorepoTest-PR

251fef6

Use REST API to find files changed instead of source checkout

dadfc61

Need to checkout workflow scripts

81d6161

Update pr-auto-label.yml

ecf10ad

github-actions Bot added the project: hipsparselt label Apr 30, 2025

add rocblas file to see auto-label in action

443dcba

jayhawk-commits requested review from NaveenElumalaiAMD, TorreZuk, amcamd, babakpst, bragadeesh, daineAMD, mahmoodw, rkamd and yoichiyoshida as code owners April 30, 2025 03:07

jayhawk-commits added 2 commits April 29, 2025 23:11

Loop through labels

2533007

Fix tabbing for env from last commit

d964751

github-actions Bot added the project: rocblas label Apr 30, 2025

Revert "[hipSPARSELt] monorepo test"

54b64c0

This reverts commit 82a24bb.

github-actions Bot removed the project: hipsparselt label Apr 30, 2025

Merge branch 'develop' into joseph-monorepoTest-PR

c3aa097

jayhawk-commits closed this Apr 30, 2025

jayhawk-commits deleted the joseph-monorepoTest-PR branch April 30, 2025 18:02

assistant-librarian Bot pushed a commit that referenced this pull request May 13, 2025

implement missing radix rank tests and missing rank_keys_desc funct…

ca394e8

…ion (#5) * added addition unit tests and implement missing function in match * updated changelog

jayhawk-commits self-assigned this May 17, 2025

jayhawk-commits added the migration Tasks or issues tied to migration to this monorepo label May 17, 2025

ozturkosu mentioned this pull request Jun 2, 2026

[CK_TILE] Add Tile Engine <-> dispatcher GEMM parity suite #7875

Closed

rrawther mentioned this pull request Jul 9, 2026

fix(rpp): Remove Ofast compiling option from CMakeList.txt and change uchar type #9237

Merged

1 task

ozturkosu mentioned this pull request Jul 15, 2026

feat(ck-tile): microscaling (mx) GEMM TE to dispatcher bridge #9329

Merged

talumbau mentioned this pull request Jul 16, 2026

feat(hipblaslt): GEKO — GEMM Kernel Optimizer #8832

Merged

tony-davis mentioned this pull request Jul 22, 2026

refactor(hipblaslt): exclude known-bug params at instantiate time #9735

Draft

1 task

ozturkosu mentioned this pull request Jul 24, 2026

feat(ck-tile): abquant GEMM TE to dispatcher bridge #9981

Closed

tony-davis mentioned this pull request Aug 4, 2026

feat(hipdnn): implement remaining cuDNN frontend graph nodes #9738

Merged

3 tasks

luraess mentioned this pull request Aug 12, 2026

Segfaults in ConvHipImplicitGemmGroup{Fwd,Wrw,Bwd}Xdlops on gfx942 (MI300) for small/degenerate convolutions — MIOpen 7.2.3 #9088

Open

Navigation Menu

Uh oh!

[hipSPARSELt] monorepo test - #5

[hipSPARSELt] monorepo test#5
jayhawk-commits wants to merge 12 commits into
developfrom
joseph-monorepoTest-PR

jayhawk-commits commented Apr 30, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

jayhawk-commits commented Apr 30, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant