[newling]

rocJITsu has specialized SIMD implementations for f16- and bf16-input MFMA
instructions with f32 accumulation. They snapshot register operands,
bulk-convert the packed inputs, perform the matrix multiplication in host SIMD
chunks, and publish staged results without per-element register access or
heap-backed result staging.

These implementations previously required native_simd<float> to have exactly
16 lanes. This enabled the fast paths in an AVX-512-targeted build, but portable
and AVX2-targeted builds fell back to generic implementations even when host
SIMD was available.

This change gives the f32-, f16-, and bf16-input f32-accumulating MFMA
specializations one shared, constexpr support predicate. The f16 and bf16
paths now use the SIMD width selected by the compiler, provided the MFMA column
count divides evenly into that width. Builds without usable SIMD, incompatible
shapes, non-wave64 execution, operand permutations, and executions with
RJ_FORCE_SCALAR=1 continue to use the existing generic paths. The
register-observation behavior used by execution plugins is preserved.

No rocJITsu-specific build option or environment variable is required to
enable the path. std::experimental::native_simd selects the width at compile
time from the compiler target. A portable x86-64 build may use four f32 lanes,
an AVX2-targeted build commonly uses eight, and an AVX-512-targeted build
commonly uses sixteen. CPU capability alone does not widen an already-compiled
binary.

fp8/bf8 MFMA, integer MFMA, and low-precision WMMA are intentionally out of
scope. CDNA3 fp8/bf8 uses FNUZ variants that need additional exact-test support,
generated i8 MFMA handlers currently use a different generic execution path,
and WMMA is a separate wave32 family with packed-output semantics.

This complements the recently merged matrix-execution improvements:

• perf(rocjitsu): use native SIMD width for f32 MFMA #10702 — use native SIMD width for f32 MFMA;
• perf(rocjitsu): snapshot scalar f32 MFMA operands #10706 — snapshot scalar f32 MFMA operands; and
• perf(rocjitsu): use native SIMD width for f32 WMMA #10710 — use native SIMD width for f32 WMMA.

Performance

Performance was measured against exact parent 79f87b823 on an AMD EPYC 9575F
host. Both revisions used AMD Clang 22, Release, -O3 -DNDEBUG, and no
explicit -march, AVX2, or AVX-512 flag. This portable build uses four-lane
native_simd<float>.

The gfx950 workload ran one simulator thread and pinned the process tree to one
logical CPU:

hipblaslt-bench --precision f16_r \
  --sizem 1024 --sizen 1024 --sizek 1024 \
  --initialization zero --iters 1 --cold_iters 0

Both binaries were warmed before four measured pairs, which alternated
execution order.

Pair	Parent wall (s)	Candidate wall (s)	Parent GEMM (s)	Candidate GEMM (s)
1	14.11	10.23	12.6251	8.61447
2	14.16	10.44	12.5957	8.92401
3	14.53	10.59	12.9857	8.90274
4	14.40	10.51	12.7722	8.88524
Median	14.28	10.475	12.6987	8.89399

The complete-process median improves by 1.36x, and hipBLASLt's reported
GEMM interval improves by 1.43x.

A two-pair bf16 1024x1024x1024 check improved complete-process wall time from
18.79 to 14.49 seconds and the reported GEMM interval from 17.08 to 12.72
seconds, or 1.30x and 1.34x, respectively.

Representative CDNA4 microbenchmarks compare the generic normal-SIMD path used
by the parent with the newly enabled specialized path:

MFMA input	Parent (ns/instruction)	Candidate (ns/instruction)	Speedup
f16	28,273	12,120	2.33x
bf16	26,843	11,360	2.36x

A separate throughput-plugin pair executed identical work in both revisions:
eight dispatches, 6,169,103 wave instructions, and 131,072 matrix
instructions.

Throughput-plugin metric	Parent	Candidate	Speedup
Active-dispatch throughput	0.4235 MIPS	0.5660 MIPS	1.34x
Matrix-handler throughput	0.01785 MIPS	0.03653 MIPS	2.05x
GEMM dispatch time	9.108 s	5.413 s	1.68x

The throughput plugin adds observer overhead, so these values are compared only
within the matched plugin-enabled pair.

Validation

• Focused MFMA, lazy-storage, and execution-plugin selection: 23 passed and one
  unrelated 16-lane-only i8 case skipped.
• All specialized f16 and bf16 MFMA shapes match the forced-scalar
  implementation bit for bit, including destructive accumulator overlap on
  the four-lane path.
• Nonzero f16 and bf16 256x256x256 hipblaslt-bench --verify runs completed
  with norm_error=0.
• Changed-file pre-commit checks passed.

Issue Tracking

Related: #10852

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.

[newling]

Hi @IanWood1, this is the last PR in the 'easy' matmul speed-ups that I have. It is essentially the same as the one for f32, but for bf16 and f16. Could you take a look please?

[accauble]

Looks like this got buried by accident? LGTM

[newling]

Thanks for the review @accauble I've done the 2 function -> 1 function suggestion you made

[newling]

Flake:
https://github.com/ROCm/rocm-systems/actions/runs/33666394067/job/100369955140?pr=10898

Error: Error: failed to run script step: command terminated with non-zero exit code: error executing command [sh -e /__w/_temp/79f95250-a6ff-11f1-b14f-bbf1b7074d59.sh], exit code 1
Error: Process completed with exit code 1.
Error: Executing the custom container implementation failed. Please contact your self hosted runner administrator.

[systems-assistant]

TheRock Submodule Bump Activity

Newest first

• 2026-09-08 — Included in TheRock via ROCm/TheRock#7974.