[CK_TILE] Add Tile Engine -> Dispatcher bridge for Stream-K GEMM

[ozturkosu]

Summary

Routes the stream_k GEMM variant through the same Tile Engine (TE) →
Dispatcher bridge already landed for regular GEMM (Phase 1, #8123) and grouped
GEMM (Phase 3, #8130). Goal of the overall effort: the Dispatcher is the single
source of truth for codegen/build/runtime, and TE only produces configs +
benchmarks.

This PR is stacked on muozturk/dispatcher-gemm-bridge (#8123) — please merge
that first. Its own diff is just two commits:

1. [CK_TILE] Add stream_k variant to GEMM Dispatcher codegen (cherry-picked)
2. [CK_TILE] Add Tile Engine -> Dispatcher bridge for Stream-K GEMM (this work)

What Stream-K needs that regular GEMM doesn't

Stream-K is a single-problem GEMM (one A/B/C, one M/N/K) with the same C
ABI as regular GEMM, so the Python side (GpuGemmRunner / GemmDispatcherLib
/ GemmProblem) and the GPU worker are reused unchanged. The differences are
internal to the .so:

• The generated launch has a Stream-K signature
  SelectedKernel::launch(const ck_tile::StreamKHostArgs&, const stream_config&),
  which allocates the reduction workspace internally (DeviceMem) and uses the
  Atomic reduction strategy.
• The registry path (generated_tile_backend.hpp::run()) hard-codes the
  single-problem GemmHostArgs launch and won't compile against a Stream-K
  SelectedKernel. So the Stream-K ctypes lib bypasses the registry and calls
  SelectedKernel::launch(args, stream) directly, reporting the name from the
  KERNEL_NAME macro (same approach grouped uses).

Changes

New

• dispatcher/bindings/ctypes/streamk_gemm_ctypes_lib.cpp — same single-problem
  C ABI (dispatcher_run_gemm(A,B,C,M,N,K,time_ms)); hipMalloc + copy A/B,
  memset C=0 (Atomic accumulates into C), build StreamKHostArgs with rcr
  strides (stride_A=K, stride_B=K, stride_C=N, k_batch=1), launch, copy C back.
  Returns 0 / -1 (HIP or throw) / -2 (kernel reports args unsupported).
• tile_engine/ops/gemm/streamk_gemm_full_benchmark.py — 3-phase driver
  (expand configs → setup_multiple_gemm_dispatchers build → subprocess-isolated
  benchmark), mirroring gemm_full_benchmark.py with variant="stream_k".
• tile_engine/ops/gemm/run_one_streamk_gemm_kernel.py — disposable GPU worker
  (identical to the regular worker since the ABI matches).
• tile_engine/ops/gemm/gemm_streamk/configs/default_config.json — small sweep
  config (128x128x{32,64}, 2x2x1, 32x32x16, compv3/compv4, intrawave, cshuffle,
  pad true, persistent false) → 4 kernels.

Modified

• dispatcher/python/gemm_utils.py — _ctypes_source_name() selects
  streamk_gemm_ctypes_lib.cpp for variant=="stream_k" (in both
  _build_compile_jobs and setup_multiple_gemm_dispatchers); .name appends
  _streamk; variant threaded into codegen_args and expand_sweep.
• dispatcher/python/ctypes_utils.py — pass the requested variant to codegen
  --variants instead of hard-coding "standard".

Validation (gfx942 / MI300X, fp16 / rcr)

Numeric parity vs a numpy fp32 reference (A.f32 @ B.f32). Stream-K's Atomic
reduction does multiple fp16 atomic-adds (one per K-split partial) vs
regular/grouped's single fp32→fp16 store, so it is inherently noisier; tolerance
is widened to max_rel ≤ 2.5e-3, frob_rel ≤ 1.5e-3 (regular/grouped use 5e-4).

Shape (M,N,K)	status	TFLOPS	max_rel	frob_rel	result
1024 x 1024 x 1024	0	23.1	1.10e-3	5.47e-4	PASS
2048 x 2048 x 2048	0	129.5	7.02e-4	3.72e-4	PASS
512 x 768 x 4096	0	24.6	2.01e-3	8.19e-4	PASS
256 x 256 x 256	0	2.1	3.44e-4	2.08e-4	PASS
4096 x 4096 x 1024	0	149.8	5.89e-4	2.65e-4	PASS

Full TE driver run (4 kernels x 4 problems = 16/16 OK, 0 failures),
default problem set uses Stream-K's sweet spot (squares + a large-K skinny shape):

• 1024 x 1024 x 1024
• 2048 x 2048 x 2048
• 4096 x 4096 x 4096
• 512 x 512 x 8192 (large-K skinny)

All status 0, positive TFLOPS, nonzero output. Name parity holds end-to-end:
the runtime name reported by each .so equals
GemmKernelConfig(variant="stream_k").name, ending in _streamk.

Unsupported-shape handling: a tiny 257^3 problem is correctly reported as
unsupported by the kernel (status -2, too few tiles to partition across CUs)
and surfaced gracefully by the bridge — not a crash.

Test plan

• unified_gemm_codegen.py ... --variants stream_k emits a *_streamk.hpp
  whose stem == GemmKernelConfig(variant="stream_k").name
• setup_multiple_gemm_dispatchers builds the Stream-K config set → .so
  compiles & links against streamk_gemm_ctypes_lib.cpp
• Numeric parity (table above) passes under the fp16 Atomic tolerance
• Full driver run 16/16 OK; name parity verified
• Unsupported tiny shape returns status -2 gracefully

Next

Land #8123, then this; afterwards delete the legacy tile_engine/ops/gemm_streamk/ machinery (Phase 4).

---

Update 2026-06-12 — brought current with #8123 + Copilot fixes

This branch had forked at #8123's first commit, so it lacked every later
regular-GEMM bridge improvement. Merged the current #8123 HEAD and ported the
Stream-K-specific analogues (the Stream-K bridge keeps its own driver, worker and
ctypes lib, so those fixes do not arrive via the merge):

• Benchmark-param parity / perf-methodology fix (streamk_gemm_ctypes_lib.cpp):
  benchmark knobs defaulted to warmup=3/repeat=10 — a cold, un-ramped clock, the
  root cause of the regular bridge's spurious "perf gap." Now default to old-TE's
  warmup=50/repeat=100, env-overridable via
  CK_TILE_BENCH_WARMUP/REPEAT/FLUSH/ROTATING. rotating_count stays 1 for
  Stream-K: the Atomic preprocess re-zeros only the original C buffer, so rotating
  C would leave rotated copies un-zeroed and corrupt the accumulation.
• --verify correctness gate (driver + worker): opt-in fp32 numpy reference
  check (global max|out-ref|/max|ref|, verified/max_rel in the CSV); a
  mismatch counts as a failure.
• Multi-GPU benchmarking (driver): fans the (kernel × problem) work across all
  visible GPUs via device-pinned HIP_VISIBLE_DEVICES workers (--devices, device
  CSV column); also fixes a latent proc-unbound error in the batch handler.
• --dtype/--layout guards (driver): constrained to the supported fp16/rcr
  surface so a mismatch fails fast.
• Copilot codegen nits: std::stoi → std::stoll for M/N/K in
  03_streamk_gemm_driver.cpp; stride-aware C zeroing via hipMemset2DAsync
  (CLayout-aware, checked HIP status) in _launch_function_streamk.
• README: new Stream-K bridge subsection.

Validation status: DONE on gfx942/MI300X (ctr-cx64-mi300x-4, enroot container). Bridge build+benchmark+--verify on gemm_streamk/configs/default_config.json = 16/16 OK, all verified (max_rel ≤ 3.0e-3, fp16 atomic tol), name parity holds. Bridge-vs-Old-TE parity (perf + correctness, byte-identical device kernel) posted as a comment with the full table + streamk_bridge_oldTE.csv.

[Copilot]

Pull request overview

This PR extends the Composable Kernel Tile Engine (TE) → Dispatcher “single source of truth” bridge to the Stream-K GEMM variant by adding Stream-K-aware codegen, a dedicated ctypes bridge that launches Stream-K kernels directly, and TE-side benchmark/worker plumbing that reuses the existing single-problem GEMM Python ABI.

Changes:

• Add Stream-K as a first-class GEMM codegen variant (--variants stream_k), including Stream-K kernel includes and a Stream-K launch wrapper that allocates/reuses internal workspace.
• Introduce a Stream-K-specific ctypes bridge (streamk_gemm_ctypes_lib.cpp) that bypasses the registry and calls SelectedKernel::launch(StreamKHostArgs, stream_config) directly while keeping the same C ABI as standard GEMM.
• Add TE benchmark driver/worker and a small default sweep config for Stream-K benchmarking via the dispatcher bridge.

Reviewed changes

Copilot reviewed 9 out of 9 changed files in this pull request and generated 2 comments.

Show a summary per file

File	Description
projects/composablekernel/tile_engine/ops/gemm/streamk_gemm_full_benchmark.py	New 3-phase TE benchmark driver for Stream-K using variant="stream_k" and subprocess isolation.
projects/composablekernel/tile_engine/ops/gemm/run_one_streamk_gemm_kernel.py	New subprocess worker to run Stream-K kernels via the existing GpuGemmRunner ABI.
projects/composablekernel/tile_engine/ops/gemm/gemm_streamk/configs/default_config.json	New default Stream-K sweep config JSON (small config set).
projects/composablekernel/dispatcher/python/gemm_utils.py	Route Stream-K builds to the Stream-K ctypes source; thread variant into codegen args and sweep expansion.
projects/composablekernel/dispatcher/python/ctypes_utils.py	Pass requested variant into unified_gemm_codegen.py --variants (instead of hard-coding standard).
projects/composablekernel/dispatcher/examples/gemm/cpp/03_streamk_gemm_driver.cpp	New minimal standalone Stream-K driver demonstrating the “dispatcher way” launch path.
projects/composablekernel/dispatcher/codegen/unified_gemm_codegen.py	Add stream_k variant: naming, includes, launch wrapper, config filtering, CLI choices.
projects/composablekernel/dispatcher/codegen/arch_filter.py	Add operator type and constraints entry for Stream-K GEMM validation.
projects/composablekernel/dispatcher/bindings/ctypes/streamk_gemm_ctypes_lib.cpp	New ctypes bridge for Stream-K that launches the force-included kernel directly and reports KERNEL_NAME.

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[ozturkosu]

Stream-K Bridge vs Old TE — MI300X parity (perf + validation)

Ran the requested comparison on ctr-cx64-mi300x-4 (AMD Instinct MI300X / gfx942, inside the enroot container, ROCm/HIP 7.13), using tile_engine/ops/gemm/gemm_streamk/configs/default_config.json → 4 kernels: {compv3,compv4} × 128x128x{32,64}, 2x2x1, 32x32x16, atomic reduction, pad-all, persistent=false, fp16 / rcr.

Shapes: 1024³, 2048³, 4096³, 512×512×8192 (the bridge's default Stream-K problem set).

Method

• Bridge: streamk_gemm_full_benchmark.py default_config.json --verify (this PR; warmup=50/repeat=100 — the new streamk_gemm_ctypes_lib.cpp defaults).
• Old TE: develop-branch standalone benchmark_gemm_streamk_* binaries (CMake trimmed to gemm+gemm_streamk only). Same warmup/repeat; perf with -verify=0, correctness with -verify=1 separately (the atomic kernel forces repeat=1 cold under verify). Verified Old TE's -flush_cache/-rotating_count change results <1% here, so the numbers are apples-to-apples.
• The generated device kernel is byte-identical between the two paths (same StreamKKernel codegen); only host-side workspace ownership / harness differs.

Validation

Both paths correct on all 16 measurements. Bridge: 16/16 verified=True, max_rel ≤ 3.0e-3 (within the fp16 atomic tolerance; Atomic does multiple fp16 atomic-adds so it is inherently noisier than a single fp32→fp16 store). Old TE: 16/16 correct.

Performance (TFLOPS, Bridge vs Old TE)

kernel	shape	Bridge	Old TE	Δ
compv3 128x128x32	1024³	24.08	23.77	+1.3%
compv3 128x128x32	2048³	136.93	137.02	−0.1%
compv3 128x128x32	4096³	260.39	257.44	+1.1%
compv3 128x128x32	512×512×8192	38.42	38.20	+0.6%
compv4 128x128x32	1024³	30.40	32.11	−5.3%
compv4 128x128x32	2048³	149.63	162.88	−8.1%
compv4 128x128x32	4096³	285.70	288.01	−0.8%
compv4 128x128x32	512×512×8192	55.04	46.90	+17.4%
compv3 128x128x64	1024³	34.09	34.26	−0.5%
compv3 128x128x64	2048³	168.62	168.10	+0.3%
compv3 128x128x64	4096³	324.34	323.57	+0.2%
compv3 128x128x64	512×512×8192	49.18	49.79	−1.2%
compv4 128x128x64	1024³	52.78	54.47	−3.1%
compv4 128x128x64	2048³	195.18	206.78	−5.6%
compv4 128x128x64	4096³	240.81	241.91	−0.5%
compv4 128x128x64	512×512×8192	81.41	66.66	+22.1%

Takeaway

• Functional parity: confirmed — identical correctness on a byte-identical device kernel.
• Performance: at parity within measurement noise. All 8 compv3 points are within ±1.3%. The compv4 points scatter wider (−8%…+22%) but center near parity; since the device kernel is byte-identical, these are harness / GPU-clock-state artifacts (standalone Old-TE benchmark process vs in-harness bridge), the same effect root-caused for regular GEMM in [CK_TILE] Add Tile Engine -> Dispatcher bridge for GEMM #8123 — not a real kernel speed difference. The largest spread is on the small/skinny + compv4 cases, which are the most clock/launch-overhead sensitive.

Raw data: streamk_bridge_oldTE.csv (16 rows; columns: kernel, M, N, K, bridge_tflops, oldte_tflops, Δ%, bridge_latency_ms, oldte_latency_ms, bridge_max_rel, bridge_verify, oldte_verify).

Note: validated on gfx942/MI300X (this comparison). The PR's own gfx942 numbers in the description above were re-confirmed: 16/16 OK, name parity, all verified.

[ozturkosu]

Stream-K reduction strategy is now a codegen axis (d326bdc283)

Previously the Stream-K codegen hard-coded the Atomic reduction strategy. This change makes reduction_strategy a first-class config axis so atomic / linear / tree reductions can each be generated, named, and selected.

Codegen (unified_gemm_codegen.py)

• reduction_strategy field on KernelConfig
• encoded in key_name() (redux_*) and KernelNaming.generate() — atomic keeps the bare _streamk suffix (name parity with the original single-strategy bridge); linear/tree are disambiguated (_streamk_linear, _streamk_tree)
• _launch_function_streamk() now emits the config's StreamKReductionStrategy (the existing reset hook already zeroes C for atomic vs the device workspace for linear/tree)
• _get_configs_for_variant(STREAM_K) iterates a new streamk_config.reduction_strategy, added to the default config as [atomic, linear, tree]

Bridge (gemm_utils.py) — GemmKernelConfig.reduction_strategy threaded into .name, to_codegen_json (per-kernel streamk_config so single-kernel codegen emits exactly the requested strategy), and expand_sweep (sweep axis). ctypes_utils.py — field on its KernelConfig for parity.

arch_specs.json intentionally unchanged — Stream-K reuses the standard warp-tile combos and arch_filter reads no Stream-K-specific keys, so adding them would be dead data.

Validation (gfx942 / MI300X)

Built + ran one kernel per strategy through the isolated worker with fp32 reference verification at 256x256x4096:

strategy	kernel	verified	max_rel
atomic	..._streamk	yes	3.1e-4
linear	..._streamk_linear	yes	3.1e-4
tree	..._streamk_tree	yes	3.1e-4

Scope

This stays within the bridge architecture. It does not add the deep dispatcher-core integration (KernelKey / KernelInstance / Problem / Dispatcher + a generic Stream-K backend + ML selectability); that remains a separate effort.

[ozturkosu]

Update: ported the Stream-K bridge onto the deep-core #8094 engine + TE parity

Reworked this PR so the Stream-K TE→Dispatcher bridge runs on the rewritten #8094 "deep-core" Stream-K engine instead of the original codegen.

What changed

• Merged the deep-core feat(ck_tile): add stream_k variant to GEMM Dispatcher codegen #8094 Stream-K codegen into the bridge (3-way merge over the shared 36eb9bdc0e ancestor; KernelKey reduction fields, KernelInstance workspace virtuals, StreamK backend, dispatcher-owned reduction workspace, registry + validation driver). The codegen now emits the struct-scope kernel type with GetWorkSpaceSize/IsSupported, keeps the 2-arg internal-workspace launch() the bridge ctypes lib calls, and adds the 3-arg dispatcher launch.
• The bridge streamk_gemm_ctypes_lib.cpp keeps launching the force-included kernel directly (no registry/dispatcher symbols), so the .so is self-contained — link no longer needs libck_tile_dispatcher.a, and the build creates build/examples itself.
• Fair-parity flags: the Stream-K .so now matches Tile Engine's gemm_streamk compile flags exactly (ground truth from a TE build's compile_commands.json): -std=c++20 -fno-offload-uniform-block -mllvm --lsr-drop-solution=1 -mllvm -enable-post-misched=0 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false --offload-compress. (-enable-post-misched=0 applied unconditionally, matching TE for streamk; -enable-noalias-to-md-conversion=0 is not used.)

Validation (fresh builds both sides, same node compiler, warmup50/repeat100, flush off, rotating 1, gpu timer, median-of-3 interleaved; verify in a separate pass):

• Codegen + on-GPU run + verify PASS for atomic / linear / tree.
• MI350 (gfx950): full parity — 12/12 cases (atomic/linear/tree × 1024³/2048³/4096³/512×512×8192) |gap%| ≤ 0.59%, all correct, no gfx950 flag changes needed.
• MI300X (gfx942): representative sweep at parity (worst real case ≈ tree −3.8%); a full sweep over the entire TE Stream-K config file (1080 fp16/rcr configs) is running, with automatic standalone re-measurement of any |gap|>15% to remove multi-GPU concurrency artifacts. Aggregate table to follow.

[ozturkosu]

Update: removed rcr hardcoding — bridge is now layout-generic (rcr/rrr/ccr/crr) + bf16

Following review, nothing layout- or dtype-specific is hardcoded anymore; strides are derived from the kernel's actual layouts everywhere.

• streamk_gemm_ctypes_lib.cpp: stride_A/B/C derived at compile time from the force-included kernel's ALayout/BLayout/CLayout (RowMajor RxC → ld=C, ColumnMajor → ld=R) instead of the old hardcoded K/K/N.
• generated_tile_backend_streamk.hpp (registry path): same layout-derived strides.
• GpuGemmRunner: reads dtype and layout off the kernel name; arranges each operand per layout (RowMajor=C-contiguous, ColumnMajor=F-contiguous); bf16 encode is memory-order-preserving so column-major operands stay column-major.
• run_one_streamk_gemm_kernel.py: dtype/layout-aware A/B + reference (was fp16-only).
• streamk_gemm_full_benchmark.py: SUPPORTED_LAYOUTS = rcr/rrr/ccr/crr, SUPPORTED_DTYPES = fp16+bf16.

Validated on GPU (verify on):

• fp16, all 4 layouts rcr/rrr/ccr/crr — all PASS, identical max_rel_err = 8.2e-4.
• bf16 + ccr (column-major A/B) — PASS, max_rel_err = 8.0e-3.
• bf16 atomic/linear/tree on MI350 — PASS.

Remaining (tracked): fp8/bf8/int8 need runner codecs (codegen + Old-TE already support them); the ctypes bridge still runs one kernel per .so (registry multi-kernel path exists in the engine via the 04 driver).

[ozturkosu]

Update: full Old-TE dtype coverage — fp8/bf8 added (commit b6bea8a)

The Stream-K bridge runner now covers the exact runnable dtype set Old-TE Stream-K supports.

Added fp8 (E4M3 FNUZ) and bf8 (E5M2 FNUZ) to GpuGemmRunner (codec ported from the plain-GEMM fp8 bridge #8887): bit-exact decode to the device fp8_t/bf8_t, nearest-representable+saturate encode, output dtype fp8/bf8 → fp16. The fp16/bf16 paths and the layout-generic (rcr/rrr/ccr/crr) logic are unchanged. SUPPORTED_DTYPES is now fp16, bf16, fp8, bf8.

GPU-validated (gfx942/MI300X, 2048³, reference = decode(encode(A)) @ decode(encode(B))):

dtype	atomic	linear	tree	col-major
fp8	PASS 192 TF, rel 9.4e-4	PASS 180	PASS 183	ccr PASS
bf8	PASS 192 TF, rel 7.8e-4	PASS 181	PASS 181	crr PASS

Exact-equivalence check vs Old-TE:

• Old-TE Stream-K's instance builder accepts fp16, fp8, bf16, bf8, fp32, fp64. fp32/fp64 have no MFMA warp tiles (not runnable on either engine). So Old-TE's runnable set = {fp16, bf16, fp8, bf8} = what the bridge now supports.
• int8 is in the warp-tile table but Old-TE's Stream-K builder rejects --datatype int8 (invalid choice), and the ck_tile engine has no warp_gemm_dispatcher specialization for int8 streamk CompV3 (compile fails). So neither TE nor the bridge generates int8 streamk. The runner keeps an int8 codec + int32-output path ready for if/when the engine adds that instantiation.

Arch guard: the fp8/bf8 codec is FNUZ (gfx942). gfx950/MI350 uses OCP fp8 (different bit layout); the runner detects arch and raises a clear error for fp8/bf8 on non-gfx942 rather than silently producing wrong results.

[ozturkosu]

Status summary — Stream-K bridge ready for review

Implementation complete and on this branch:

• Ported onto the deep-core feat(ck_tile): add stream_k variant to GEMM Dispatcher codegen #8094 engine (KernelKey reduction fields, workspace virtuals, StreamK backend, dispatcher-owned reduction workspace, registry/validation driver) via a clean 3-way merge.
• Layout-generic — rcr/rrr/ccr/crr; strides derived from the kernel's ALayout/BLayout/CLayout (no rcr hardcoding) in both the ctypes lib and the registry backend.
• Full Old-TE dtype coverage — fp16, bf16, fp8 (E4M3 FNUZ), bf8 (E5M2 FNUZ). This is the exact runnable set Old-TE Stream-K supports: TE's instance builder also has only {fp16, fp8, bf16, bf8} runnable (fp32/fp64 have no warp tiles; int8 is rejected by TE's builder and has no ck_tile warp_gemm streamk specialization). int8 codec is kept ready for if/when the engine adds it.
• Fair flags — the .so matches TE's gemm_streamk compile flags exactly (verified against a real compile_commands.json).
• Reductions — atomic / linear / tree, all verified.

Verification (verify on): fp16 all 4 layouts PASS (rel 8.2e-4); bf16 + col-major PASS; fp8/bf8 × atomic/linear/tree × row+col-major PASS (rel <1e-3, gfx942).

TE↔bridge parity (fair, fresh both sides, median-of-3 interleaved):

• MI350 (gfx950): full pass — worst |gap| 0.59%.
• MI300X (gfx942): full canonical-config sweep in progress — ~89% within ±15%, median |gap| ≈ 2.75%; the >15% outliers are small-shape multi-GPU concurrency artifacts, auto re-measured standalone.

Caveat: fp8/bf8 codec is gfx942 FNUZ; the runner raises a clear error on gfx950 (OCP) pending separate handling.

[ozturkosu]

Clarification: fp8/bf8 Stream-K on gfx950 (MI350) — blocked by a shared upstream ck_tile bug, not the bridge

For completeness on why MI350 dtype coverage is scoped to fp16+bf16:

A pre-existing, shared ck_tile engine bug on gfx950. fp8/bf8 Stream-K on gfx950 produces output ≈ 4× the reference (e.g. 2792 vs 698). Verified facts:

• fp16/bf16 Stream-K verify correct on gfx950; only fp8/bf8 fail.
• Old-TE fails with the identical 4× error — both engines instantiate the same StreamKKernel template, so this is a pre-existing shared ck_tile kernel bug, not introduced by this bridge/PR.
• gfx942 (MI300X) fp8/bf8 verify fine on both engines → the bug is gfx950-specific.
• The 4× factor == the K-split partial-accumulation factor → root cause is in the fp8/bf8 reduction/accumulation path on CDNA4. Needs an upstream ck_tile fix.

Bottom line: even with an OCP fp8 codec (gfx950 uses OCP vs gfx942 FNUZ), the gfx950 fp8/bf8 Stream-K kernels themselves are broken — and Old-TE is equally broken — so there is nothing correct to be "at parity" with. fp8/bf8 parity is therefore only meaningful on gfx942/MI300X, where both engines work and the bridge is at parity. On MI350 the valid dtypes are fp16 + bf16 (both swept).

That's why MI350 is scoped to fp16/bf16, and gfx950 fp8/bf8 is flagged as a separate upstream ck_tile issue rather than a bridge gap. (The bridge runner guards this: fp8/bf8 raise a clear error on any non-gfx942 arch instead of silently producing wrong results.)

[ozturkosu]

Parity analysis (interim — full canonical-config sweeps)

Running the full TE Stream-K config matrix (every bridge-generatable tile/pipeline/persistent × atomic/linear/tree × 4 shapes) per arch/dtype, fair flags (warmup50/repeat100, flush off, rotating 1, gpu timer, median-of-3 interleaved per case, both sides built fresh on the same node compiler), with standalone auto-recheck of any |gap|>15%.

Aggregate so far (gap% = (bridge−TE)/TE; aggregates over measured cases):

GPU	arch	dtype	measured	within ±15%	median |gap|	max |gap|
MI300X	gfx942	fp16	2090	93%	1.42%	28.2%
MI350	gfx950	bf16	1464	100%	0.63%	7.7%

Analysis:

• bf16 on MI350 (single GPU) is 100% within ±15%, median 0.63%, zero outliers — i.e. with no GPU contention the bridge matches TE essentially exactly.
• fp16 on MI300X is 93% within ±15%, median 1.42%. The ~7% over ±15% are small/sub-millisecond shapes (tiny 64×64 tiles) measured under 4-GPU concurrency. Standalone re-measure confirms these are the TE standalone-benchmark per-launch overhead on tiny kernels (the bridge runs faster there because the .so path has less harness overhead) — the device kernels are identical, so it's a measurement artifact, not a codegen regression. They collapse toward parity when re-measured single-GPU (the contention-free MI350 run shows 0 such outliers).
• Reduction strategies atomic / linear / tree all covered and at parity.

Still running (final tables to follow): MI300X bf16/fp8/bf8 (gfx942) and MI350 fp16 (gfx950), in parallel across allocations. Raw per-row results (~5K) are kept in CSVs; only this analysis is posted here.

Bottom line so far: the Stream-K bridge is at parity with Tile-Engine across the swept configs; deviations are measurement artifacts on tiny kernels, not real gaps.

[ozturkosu]

Result CSV files (raw per-row data)

All raw parity results live on the build host under /home/AMD/muozturk/ (these will be consolidated into /home/AMD/muozturk/PR8136_parity_csvs/ once all sweeps finish). Per-arch/per-dtype sheets + one combined master:

Master (all archs+dtypes, with gpu + dtype label columns):

• PR8136_streamk_te_bridge_parity_ALL.csv

Per arch / dtype:

• streamk_bridge_te_parity_MI300X_gfx942_FULL.csv (MI300X, fp16)
• streamk_bridge_te_parity_MI300X_gfx942_bf16_FULL.csv (MI300X, bf16)
• streamk_bridge_te_parity_MI300X_gfx942_fp8_FULL.csv (MI300X, fp8)
• streamk_bridge_te_parity_MI300X_gfx942_bf8_FULL.csv (MI300X, bf8)
• streamk_bridge_te_parity_MI350_gfx950_bf16_FULL.csv (MI350, bf16)
• streamk_bridge_te_parity_MI350_gfx950_fp16_FULL.csv (MI350, fp16)

Columns: gpu, dtype, gfx_arch, strategy, tile, wave, warp_tile, pipeline, persistent, M, N, K, te_tflops, bridge_tflops, gap_pct, bridge_max_rel_err, note (gap_pct = (bridge−TE)/TE·100; +ve = bridge faster; note flags standalone_recheck / hang / te_timeout).

(The Confluence page carries the analysis/summary only; these CSVs hold the ~5K raw rows.)

[ozturkosu]

Superseded by #9028, same commit (b6bea8a) on a policy-compliant branch (users/muozturk/ck-tile/dispatcher-te-bridge-streamk-gemm) targeting develop. Please continue review there.