Repository navigation
[rocjitsu] Fix zero-selector FLOAT32 buffer stores and enable guitar-hero coverage - #13988
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
gfx1250 B0-to-A0 translation baselineWarning The PR result is not directly comparable to the latest develop baseline. Changed provenance:
Compared PR head |
FYI, this was by design, I added test cases that covered bugs that I had already fixed locally (either in-review or blocked by other work)... |
ec5c16a to
ece83e6
Compare
ece83e6 to
cf9f4a1
Compare
cf9f4a1 to
f7990da
Compare
kuhar
left a comment
There was a problem hiding this comment.
Added 2 comments and an inline reply
f7990da to
77a7ea7
Compare
d7a8411 to
83a74c0
Compare
93e8c27 to
8e29340
Compare
8e29340 to
c0f37a4
Compare
The corpus case rocblas_sgemm.large_2048x2048 on gfx1201 runs for about 141 s in the Clang ASan+UBSan lane without vptr checks and 360-457 s with them, past the lane's 360 s first-pass timeout. run-corpus-tests.sh gains --ubsan, which deselects the cases listed in ubsan_excluded_cases from every simulator-suite pass. The workflow passes it for both UBSan lanes. Those lanes still run the smaller rocBLAS SGEMM cases, and Release still runs the large one.
rocjitsu post-commit fuzz passedThe 32-worker multiplexer campaigns for commit
|
TheRock Submodule Bump ActivityNewest first
|
Untyped buffer stores with a single-channel FLOAT32 format and all four descriptor DST_SEL fields set to zero are rejected by the emulator. The guitar-hero instruction corpus intentionally exercises this and other known emulator bugs, including fixes already under review when the cases were added. Enabling the newer corpus exposed the CI failures tracked in #12583. Native gfx1100 and gfx1201 accept this descriptor and route it like the A8 selector view.
This PR resolves the qualified zero-selector view to A8 before packing the store. W supplies the memory channel when present. When W is absent, the existing missing-channel policy supplies X on gfx1201 and zero on gfx1100. These measurements cover gfx1100 and gfx1201; the routing and zero-fill policy on the other supported targets are extrapolated. Both buffer entry points use one helper for qualification and selector resolution. Shared image packing keeps its existing behavior.
Regressions use distinct X/Y/Z/W values, including packed D16 components, and cover all nine ordinary/D16 store forms with sparse EXEC. Direct packing covers both missing-channel policies. Boundary checks through both buffer APIs reject zero selectors for single-channel integer, narrower float and multichannel formats without changing output bytes or issuing a store. Integer and multichannel image-store regressions preserve the image behavior.
The other relevant emulator fixes have landed: #13705 supplies typed integer minmax lowering, floating-dot DPP availability and per-source FP8/BF8 selection; #13691 supplies the DOT arithmetic and accumulator modifiers. This PR retains both native DOT4 fixture sets and decoded coverage of arithmetic, accumulator modifiers, plain/DPP operand handling, wave sizes, host rounding modes, sparse EXEC and overlapping registers. Related generator coverage is consolidated, and the DOT arithmetic and DPP availability contracts are documented.
Both workflow corpus references advance from e191fc2 to ebb4004d, enabling the guitar-hero cases. Bounded apt retries handle transient dependency-download failures. The separate VMEM/GDS completion work and newer corpus pin remain in #12583.
Validation: the rebased Clang Release and ASan/UBSan builds pass the affected buffer/image, DOT/WMMA and integer regressions, including the captured arithmetic cases. The affected Python suites and changed-file hooks pass. All ten guitar-hero cases pass with the exact pinned corpus and the workflow's ROCm SDK version.
Appendix: omitting the large SGEMM corpus case from UBSan lanes
The corpus case
rocblas_sgemm.large_2048x2048on gfx1201 takes about 141 s in the Clang ASan+UBSan lane without vptr checks and 360-457 s with them, which exceeds the lane's 360 s first-pass timeout. On the parent commit, the asan-ubsan lane failed on this case: its gfx1201 pass took 374 s and failed when this case timed out, and the case passed only in a 458 s rerun. The final commit adds--ubsantorun-corpus-tests.sh, which deselects the cases listed inubsan_excluded_casesfrom the simulator suites, and the workflow passes it for both UBSan lanes. Those lanes still run the smaller rocBLAS SGEMM cases, and Release still runs the large one.To check that nothing else was dropped, the table compares the corpus test summaries (passed / skipped) of run 41003966356 on this commit with run 40995568287 on its parent. The counts are identical in the ubsan-gcc and asan-ubsan lanes, apart from the parent's asan-ubsan timeout described above.
Exactly one item is removed per architecture. On gfx1201 it is a passing test; on the other four architectures the case was already skipped, so one skip disappears. No other passed count changes, and the case name no longer appears in either UBSan log. The native ctest counts are unchanged, with all 9845 tests passing in ubsan-gcc and all 9857 in asan-ubsan. After the change, the gfx1201 pass takes 44 s in asan-ubsan instead of 374 s. The release lane still passes the large case on gfx1201, and the tsan lane also passes. For comparison, the parent run's ubsan-gcc job passed the large case in its gfx1201 pass, which took 114 s.