[CK_TILE] TE->Dispatcher GEMM bridge (consolidated: key + bf16 + layouts + parity + example)#8261
Closed
ozturkosu wants to merge 96 commits into
Closed
[CK_TILE] TE->Dispatcher GEMM bridge (consolidated: key + bf16 + layouts + parity + example)#8261ozturkosu wants to merge 96 commits into
ozturkosu wants to merge 96 commits into
Conversation
Introduce the shared-config bridge that lets Tile Engine drive the GEMM dispatcher the same way FMHA and grouped conv already do: one config dataclass owned by the dispatcher, imported by Tile Engine, with no translator between two vocabularies. - dispatcher/python/gemm_utils.py: GemmKernelConfig (the common contract; .name mirrors the codegen KERNEL_NAME byte-for-byte), GemmProblem, GemmDispatcherLib, GpuGemmRunner, setup_multiple_gemm_dispatchers (codegen + hipcc -> .so paths, CPU-only/parallel), and expand_sweep. - dispatcher/bindings/ctypes/gemm_ctypes_lib.cpp: add the indexed multi-kernel ABI dispatcher_get_kernel_name_at(index, buf, size); legacy single-kernel dispatcher_get_kernel_name retained. - tile_engine/ops/gemm/gemm_full_benchmark.py: 3-phase TE driver (compile -> load problems -> benchmark) that generates no binaries. - tile_engine/ops/gemm/run_one_gemm_kernel.py: disposable GPU worker for subprocess fault isolation. Scope: regular GEMM, fp16, rcr (Phase 1). Name parity verified end-to-end (config.name == generated .hpp stem == runtime registry name). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The regular GEMM ctypes bridge (dispatcher_initialize) registered every force-included kernel under a hard-coded KernelKey (fp16 / rcr / 128x128x32 / compv4-cshuffle-intrawave). Any kernel whose real traits differed -- e.g. a compv3, 64x64x64, or default-epilogue specialization -- was registered under a key that lied about it. This is the core generality gap vs the FMHA reference (#5260), which enumerates real configs. Codegen now emits GEMM_KEY_* macros into the CK_TILE_SINGLE_KERNEL_INCLUDE block describing the kernel's actual dtypes, layouts, tile/wave/warp_tile shapes, pipeline, scheduler, epilogue, block size and flags. The ctypes library builds the KernelKey from those macros using the existing string_to_* helpers in kernel_key.hpp (their accepted strings match the codegen's emitted values byte-for-byte). The legacy hard-coded block is kept behind #else as a default for headers generated before the macros existed. Validated on gfx942/MI300X: macros match each kernel struct's own constexpr values; isolated key-build asserts pass; full TE bridge run (compv3/cshuffle + compv3/default + persistent variants x 512/1024) = 6/6 OK, positive TFLOPS, name parity intact. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Closes the regular-path dtype-coverage gap (fp16-only) identified when measuring the GEMM bridge against the FMHA reference (PR #5260). The C ABI is already dtype-agnostic (sizes buffers from sizeof(ADataType)) and the codegen already emits bf16 kernels; the only fp16 hard-coding lived in the Python host buffers. - gemm_utils.py: GpuGemmRunner.run() is now dtype-aware. It detects the kernel's real dtype from the compiled kernel name and encodes the host buffers to match. bf16 has no native numpy dtype, so it is carried as a uint16 bit pattern (sizeof(bf16_t) == 2): _fp32_to_bf16_u16 encodes fp32 -> bf16 with round-to-nearest-even, _bf16_u16_to_fp32 decodes the output back to fp32. fp16 path is unchanged. - run_one_gemm_kernel.py: worker now generates fp32 source data so the runner owns all dtype encoding (no double rounding). Stacked on muozturk/dispatcher-gemm-registry-key (PR #8187): bf16 needs the trait-derived KernelKey so the registry reports the correct dtype. Validated on gfx942 (MI300X): 3 bf16 kernels x 2 problems = 6/6 OK, 46-219 TFLOPS; numeric check vs CPU reference max_rel_err 7.7e-3 (within bf16 tolerance). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Closes the regular-path layout-coverage gap (rcr-only) vs the FMHA reference (PR #5260). The codegen, flat C ABI, and TE driver already supported arbitrary A/B/C layouts (expand_sweep threads --layout into layout_a/b/c; the kernel computes strides from its compiled layout + M,N,K). The only rcr-locked spot was the host-buffer transpose in GpuGemmRunner.run(), which hard-coded B.T. - gemm_utils.py: run() is now layout-aware. It parses the 3-char layout token from the compiled kernel name (gemm_<dtype>_<layout>_...) and lays out the host buffers to match: a 'c' (column-major) operand is transposed so the contiguous buffer's flat memory equals column-major order; C is allocated (M,N) for row-major or (N,M) for column-major and transposed back to logical MxN. Composes with the fp16/bf16 dtype handling. - _layout_from_kernel_name() helper added. Validated on gfx942 (MI300X), M=N=K=512, numeric check vs CPU reference: fp16 rcr/rrr/ccr/crr -> PASS (max_rel 5.8e-4) bf16 rrr -> PASS (max_rel 7.8e-3, dtype x layout compose) Column-major C (e.g. rcc) is rejected at build by ck_tile's universal GEMM (row-major C only) -- a kernel limitation, not a host issue. Stacked on muozturk/dispatcher-gemm-bf16 (PR #8190). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Examples 01-11 drive the Dispatcher's native ctypes Registry; none exercise the Tile Engine -> Dispatcher bridge (dispatcher/python/gemm_utils.py). Add 12_te_bridge.py, which builds via setup_multiple_gemm_dispatchers and runs via GpuGemmRunner across fp16/bf16 and the row/col layout combinations the universal GEMM supports (row-major C), validating each against a NumPy reference. Closes the last FMHA-parity gap (runnable example scripts). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two fixes so the trait-derived KernelKey stays truthful (addresses Copilot
review on the registry-key change):
- gemm_ctypes_lib.cpp: the GEMM_KEY_PAD_M/N/K macros were emitted by codegen
but never assigned to key.algorithm.pad_m/n/k. Since pad flags feed both the
key's hash/equality and the kernel name, a padding-disabled kernel registered
under a key claiming pad=true, disagreeing with its own name. Assign them.
- kernel_key.hpp: string_to_scheduler("default") silently fell through to
Intrawave. Preshuffle kernels emit "default", which codegen maps to
Scheduler::Auto (codegen_common.py SCHEDULER_TO_DISPATCHER); mirror that.
Validated on gfx942: pad-disabled fp16/rcr kernel builds, registers, and runs
(max_rel 2.3e-4 vs NumPy).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…spatcher-gemm-bf16
…her-gemm-bridge-example
Addresses the Copilot review requests on the bf16 (#8190) and layout (#8191) changes: the bit-level bf16 round-to-nearest-even encoder and the dtype/layout name parsers had no coverage. Adds dispatcher/tests/test_gemm_utils.py (10 CPU-only tests): bf16 exact round-trip, <=2^-8 relative error, tie-to-even, Inf/NaN, uint16 size; dtype/layout parsing with fallbacks; and a GemmKernelConfig.name -> parser round-trip locking the config/codegen/runtime naming contract. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…her-gemm-bridge-example
setup_multiple_gemm_dispatchers built its per-config codegen args without the config's variant, and _generate_single_kernel_subprocess hard-coded --variants standard. A GemmKernelConfig(variant='preshuffle'/'multi_d') passed to the bridge would therefore be code-generated as a standard kernel, while its name (and the hpp_glob_pattern derived from it) still carried the variant suffix -- so the lookup could never match the emitted header. Pass the variant through; the subprocess default stays "standard" so all existing callers are unaffected. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…er-gemm-registry-key
…spatcher-gemm-bf16
…her-gemm-bridge-example
…dler If subprocess.Popen() itself raises, the generic except handler referenced proc before it was bound, masking the real error with an UnboundLocalError. Initialize proc = None before the try; the handler already guards with `if proc and proc.poll() is None`. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…er-gemm-registry-key
…spatcher-gemm-bf16
…her-gemm-bridge-example
- Add --arch flag (defaults to rocminfo auto-detection) and thread it into GemmKernelConfig so the example is not pinned to gfx942. - Complete the bf16 layout matrix (add ccr/crr) so _CASES matches the docstring's claim of fp16+bf16 across all four row-major-C combinations. - Round the CPU reference's output to dtype_c as well, matching the GPU's back-store of C, so accuracy comparison is faithful end to end. Validated 8/8 on gfx942 (fp16 max_rel 4.9e-4, bf16 3.9e-3). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
_fp32_to_bf16_u16() already forces a contiguous float32 buffer internally, so wrapping its argument in np.ascontiguousarray only added an extra copy of the (possibly transposed) operand. Pass the laid-out array directly. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…her-gemm-bridge-example
Copilot review item: the driver exposed unrestricted --dtype/--layout while the Phase-1 worker hard-codes fp16 inputs and an rcr (column-major B) host transpose. Passing e.g. --dtype bf16 would codegen bf16 kernels but feed them fp16 data, silently benchmarking the wrong thing. Add SUPPORTED_DTYPES/ SUPPORTED_LAYOUTS and wire them into argparse choices so a mismatch fails fast. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…er-gemm-registry-key
…spatcher-gemm-bf16
This branch's runner encodes bf16 from the kernel name, so add bf16 to SUPPORTED_DTYPES. Layout stays rcr-only here. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This branch adds layout-aware host transpose, so accept all four row-major-C layouts. Column-major C is rejected by ck_tile at build, so it stays out. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
- ab_same_harness.py: derive ROOT from __file__ and take old-TE header dir from OLD_TE_GEN env var (was hard-coded dev paths); drop unused statistics import - generated_tile_backend.hpp: make env_bool case-insensitive (handles False/Off) and align its comment - gemm_full_benchmark.py: clarify in --devices help and resolve_devices docstring that a bare digit is a count; a single id needs the comma form (5,) Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The bridge compiled GEMM kernels with only -mllvm -enable-noalias-to-md-conversion=0, omitting the AMDGPU inlining/codegen flags Tile Engine passes (-amdgpu-early-inline-all=true, -amdgpu-function-calls=false, --lsr-drop-solution=1, -enable-post-misched=0, -fno-offload-uniform-block). Those flags change register allocation and thus occupancy. Persistent GEMM kernels size their grid by occupancy (UniversalGemmKernel::MaxOccupancyGridSize = #CUs x occupancy), so the mismatch produced large perf gaps vs Tile Engine on persistent tiles (up to ~-73% on some mem-pipeline 256x256x64 / 192x64x128 kernels) while non-persistent kernels (fixed problem-sized grid) were unaffected. Matching the flags restores parity: on the full persistent buildable-common fp16/rcr subset (596 kernels x 5 shapes), the bridge-vs-Tile-Engine gap collapses to a median ~-2% with no kernel beyond +/-20% (was -51.7% .. +203.8%). Numerically unchanged (max_rel ~3.5e-4). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
…ridge (#8261) Brings #8261 current with #8123's latest work: arch-validated tile filtering, matching Tile Engine AMDGPU -mllvm codegen flags, the --verify fp32-reference gate, perf-gap diagnosis correction, Copilot nits, and the develop catch-up. Conflict resolution (tile_engine/ops/gemm): - Adopted #8123's flat op-root configs/ layout for gemm_universal (gemm_universal/ removed; matches the fmha/grouped_conv bridge convention and the on-disk rename). - Kept #8261's superset support surface: SUPPORTED_DTYPES=(fp16,bf16), SUPPORTED_LAYOUTS=(rcr,rrr,crr,ccr) in driver + README, since #8261 already wires bf16 and all four layouts through the dispatcher host path. - Took #8123's --verify driver/worker wiring and device-count semantics. - example_problems.json: identical content, kept at configs/.
github-actions
Bot
added
project: miopen
project: hipblas
project: hipblaslt
project: hipcub
project: hipfft
project: hiprand
project: hipsparse
project: hipsparselt
project: rocblas
project: rocfft
project: rocprim
project: rocrand
project: rocthrust
shared: rocroller
github actions
shared: origami
project: hipdnn
project: hiptensor
project: miopen-provider
ci:hipsparselt-fast
project: hipdnn-integration-tests
shared: stinkytofu
labels
Contributor
Author
Synced with #8123 — latest commit
|
5 tasks
Contributor
Author
Superseded by #8479 — work moved thereClosing this PR. The GEMM bridge work has been moved to #8479 This branch's history had accumulated unrelated cross-project commits through All functionality from this PR is preserved in #8479: the TE → Dispatcher regular Please review #8479 instead. |
5 tasks
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
20 participants
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
Consolidates the full Tile Engine → Dispatcher GEMM bridge into one PR, gathering the changes previously split across the stack #8123 → #8187 → #8190 → #8191 → #8193:
GemmKernelConfigwhose.nameequals the codegenKERNEL_NAMEbyte-for-byte; flat extern "C" ABI;GpuGemmRunner.GEMM_KEY_*macros replace the hard-coded fp16/rcrKernelKey.rrr/rcr/crr/ccrvia layout-aware host transpose (column-major C is rejected at build by ck_tile's universal GEMM — kernel limitation).tests/test_gemm_parity.py(GPU-gated, NumPy oracle) andexamples/gemm/python/12_te_bridge.py(matrix / shapes / sweep gallery).Combined diff vs
develop: 10 files, +2034 / −3, all underdispatcher/andtile_engine/.Validation (gfx942 / MI300X)
Parity is checked as a global relative error
max|gpu − ref| / max|ref|against an independent NumPy reference (not a bridge-vs-TE self-comparison). A bridge-vs-native-TE head-to-head on thercrsurface confirms the bridge reproduces the Tile Engine where both are correct (bridge 4.6e-4 / 3.3e-4 vs TE verdictcorrect, threshold 4.88e-4).Current parity matrix — 10/24:
Two known defects the suite surfaces (square-only validation masked them):
rrr/ccr/crrfall back to the square-coincident stride, so they are correct only when the mismatched dims are equal.rcris correct on every shape. In the dispatcher/ABI stride plumbing, not the Python runner or test.awkward(N=129) returns status -1 for every layout incl.rcr; K=512 is ÷8 so not the K÷8 rule — the odd-N vectorized C-store. Independent of Bug 1.The square surface (8/8) and the
rcrsurface across shapes are solid; the non-native non-square cases are not yet correct.Test plan
python3 -m pytest dispatcher/tests/test_gemm_utils.py(host helpers, CPU-only)python3 dispatcher/tests/test_gemm_parity.pyon a gfx942 box (GPU-gated; skips cleanly without GPU/hipcc)python3 dispatcher/examples/gemm/python/12_te_bridge.py --demo matrix