Repository navigation
feat(dispatcher/heuristics): training data from the gemm driver's CSV - #13431
Merged
Merged
Conversation
Base automatically changed from
users/andriy/ck/2499-bf16-perf-model
to
develop
October 10, 2026 03:55
✅ All Checks Passed — Ready for Review
📖 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. |
This was referenced Oct 14, 2026
Merged
andriy-ca
marked this pull request as ready for review
October 14, 2026 19:44
andriy-ca
force-pushed
the
users/andriy/ck/heuristics-gemm-driver-csv
branch
from
October 14, 2026 20:03
212c575 to
efcf478
Compare
illsilin
pushed a commit
that referenced
this pull request
Oct 28, 2026
…model, and a runbook (#13432) JIRA ID : AICK-2523 > **Stacked on #13431.** Only the last commit is this PR. It also needs #13430's driver options at run time. ## Why A full sweep builds every kernel of a configuration space and times it on every problem it can run, which can take days for a large space over many problems. A model trained on earlier measurements can rank each problem's candidates so that only its top K are built and timed, and a sweep of the same problems measures how good the picks are. This PR adds the selection step and a runbook for the whole workflow. ## What changes - `dispatcher/heuristics/select_kernels.py`: ranks the candidates from `gemm_full_benchmark.py --list-pairs` with `Predictor.rank_kernels` and writes the top K per problem as a `--pairs` file for the driver, plus a CSV of the picks (problem, rank, kernel, predicted TFLOP/s). - The candidates are the driver's own pairs, after its vector-width and tile rules, so every pick can launch. - Ties in the predicted value break on the kernel name. - The driver's field names map to the feature engine's (`wave_*` -> `warp_*`, `vector_size_*` -> `vec_*`). - It refuses a model trained for another dtype or arch, and non-universal variants. - `dispatcher/heuristics/HEURISTIC_BENCHMARKING.md` ([view](https://github.com/ROCm/rocm-libraries/blob/users/andriy/ck/heuristics-select-kernels/projects/composablekernel/dispatcher/heuristics/HEURISTIC_BENCHMARKING.md)), a runbook: when the workflow applies, prerequisites and inputs, the reference sweep, training data and model, candidates and selection, measuring the picks, scoring them against the sweep, cost, and troubleshooting. It is linked from the heuristics README and the tile_engine GEMM README, and whitelisted in `dispatcher/heuristics/.gitignore` like the other tracked docs. ## Testing - `tests/test_select_kernels.py` (13 tests): - field mapping, top-K and tie-breaking, and the model/run checks; - an end-to-end run on the in-tree `gemm_universal_fp16_gfx950` model; - the runbook's scoring snippet, executed on constructed data; - every option the runbook's commands use exists in that script's argparse. The driver part skips until #13430 is in the tree. - On a tree merging this branch with #13430: - heuristics suite: 545 passed, 1 skipped (a pre-existing skip); the GEMM driver suites pass; - the runbook's steps run as written. Training data, `train.py`, `--list-pairs` and `select_kernels.py` ran for real; the sweep and the picks run used stand-ins for the compiler and GPU. Only the picks' kernels were built, the measured pairs equalled the picks, and the scoring snippet ran on the outputs. ## Merge order ```text #13431 -> this PR #13430: independent of #13431; must merge before this PR ``` 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
convert_csv_to_parquet.py read gemm_full_benchmark.py CSVs with its generic parser, which matches only "gemm_universal_" names: the driver's "gemm_<dtype>_" kernels kept dtype and pipeline only, without tile, warp, padding, epilogue or vector widths, under non-canonical column names, and rows that failed --verify were kept as labels. - data_pipeline.parse_gemm_driver_csv: canonical rows (CANONICAL_COLUMNS) from a driver CSV. Kernel fields, dtype and layout come from parse_kernel_name; rows with verified=False are dropped and counted; a name that does not decode raises; floats are read round-trip; bandwidth counts A and B in the input dtype and C in the driver's output dtype. - data_pipeline.main accepts a .csv input. - convert_csv_to_parquet.py detects a driver CSV, delegates to it and adds the arch's hardware constants (HW_PROFILES). Other CSVs keep the generic path. - README: the conversion step. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…model, and a runbook (#13432) JIRA ID : AICK-2523 > **Stacked on #13431.** Only the last commit is this PR. It also needs #13430's driver options at run time. ## Why A full sweep builds every kernel of a configuration space and times it on every problem it can run, which can take days for a large space over many problems. A model trained on earlier measurements can rank each problem's candidates so that only its top K are built and timed, and a sweep of the same problems measures how good the picks are. This PR adds the selection step and a runbook for the whole workflow. ## What changes - `dispatcher/heuristics/select_kernels.py`: ranks the candidates from `gemm_full_benchmark.py --list-pairs` with `Predictor.rank_kernels` and writes the top K per problem as a `--pairs` file for the driver, plus a CSV of the picks (problem, rank, kernel, predicted TFLOP/s). - The candidates are the driver's own pairs, after its vector-width and tile rules, so every pick can launch. - Ties in the predicted value break on the kernel name. - The driver's field names map to the feature engine's (`wave_*` -> `warp_*`, `vector_size_*` -> `vec_*`). - It refuses a model trained for another dtype or arch, and non-universal variants. - `dispatcher/heuristics/HEURISTIC_BENCHMARKING.md` ([view](https://github.com/ROCm/rocm-libraries/blob/users/andriy/ck/heuristics-select-kernels/projects/composablekernel/dispatcher/heuristics/HEURISTIC_BENCHMARKING.md)), a runbook: when the workflow applies, prerequisites and inputs, the reference sweep, training data and model, candidates and selection, measuring the picks, scoring them against the sweep, cost, and troubleshooting. It is linked from the heuristics README and the tile_engine GEMM README, and whitelisted in `dispatcher/heuristics/.gitignore` like the other tracked docs. ## Testing - `tests/test_select_kernels.py` (13 tests): - field mapping, top-K and tie-breaking, and the model/run checks; - an end-to-end run on the in-tree `gemm_universal_fp16_gfx950` model; - the runbook's scoring snippet, executed on constructed data; - every option the runbook's commands use exists in that script's argparse. The driver part skips until #13430 is in the tree. - On a tree merging this branch with #13430: - heuristics suite: 545 passed, 1 skipped (a pre-existing skip); the GEMM driver suites pass; - the runbook's steps run as written. Training data, `train.py`, `--list-pairs` and `select_kernels.py` ran for real; the sweep and the picks run used stand-ins for the compiler and GPU. Only the picks' kernels were built, the measured pairs equalled the picks, and the scoring snippet ran on the outputs. ## Merge order ```text #13431 -> this PR #13430: independent of #13431; must merge before this PR ``` 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
andriy-ca
force-pushed
the
users/andriy/ck/heuristics-gemm-driver-csv
branch
from
October 28, 2026 18:09
7d7e40e to
1b13188
Compare
Contributor
Author
|
@illsilin thanks for approving #13432. It was stacked on this PR, so merging it landed its commit on this branch rather than on develop: this PR now carries both the training-data step and the kernel selection + runbook from #13432. I rebased it onto current develop (one conflict with the batched GEMM path in The runbook documents the driver options in #13430, so #13430 should merge before this PR. |
illsilin
pushed a commit
that referenced
this pull request
Oct 29, 2026
) JIRA ID : AICK-2523 ## Why The tile_engine GEMM driver measures every built kernel on every problem it can run, timing each pair once. That makes three things hard: measuring only a chosen set of `(problem, kernel)` pairs (the kernels a trained model ranks highest, or the problems an interrupted sweep did not reach), trusting which kernel is fastest when it rests on single timings, and knowing why a pair is missing from the results. This PR adds the driver side of a workflow in which a model picks the kernels to measure and a sweep judges the picks; the workflow is written up as a runbook in #13431 ([HEURISTIC_BENCHMARKING.md](https://github.com/ROCm/rocm-libraries/blob/users/andriy/ck/heuristics-gemm-driver-csv/projects/composablekernel/dispatcher/heuristics/HEURISTIC_BENCHMARKING.md)). ## What changes `tile_engine/ops/gemm/gemm_full_benchmark.py` gains five options, all off by default. The helpers the other sweep drivers import (`_run_batch_on_device`, `benchmark_on_devices`) keep their signatures; their new parameters are keyword-only with defaults. - `--build-timeout SECONDS`: one limit for each kernel's codegen, compile and link step, passed through `setup_multiple_gemm_dispatchers(build_timeout=...)` (keyword-only). The codegen and link limits were hardcoded at 300 s in `ctypes_utils`; both now read a value from their job, defaulting to 300 s, and a timed-out codegen, compile or link step reports the step and the limit (`Timeout: link exceeded 300 s`) instead of `Timeout` or the subprocess exception text. - `--failures-csv PATH`: one row per failed measurement (`error`, `missing`, `worker-exit`, `timeout`, `batch-error`, `mismatch`, `bad-output`), one `unsupported` row per kernel the worker reports as not supporting the problem (counted as skipped, as before), and one `build-failed` row per failed build. `_run_batch_on_device` takes an optional `failures` list; each failed build's row carries the failed step and its first error line, from `setup_multiple_gemm_dispatchers`' `errors`. Failed builds are recorded even when no kernel builds. - `--repeat-top K`, `--repeat-count N`, `--repeat-csv PATH`: after the sweep, re-time the K fastest kernels of each problem N more times, one subprocess per launch, rep-major. Kernels that failed `--verify` or wrote no non-zero output are skipped. - `--list-pairs PATH`: dry run. Writes every pair the run would measure, with each kernel's configuration, and exits without building. - `--pairs PATH`: measure only listed pairs (`"MxNxK"` to kernel names, or `"*"`; a `--list-pairs` file is accepted). Expansion still uses the full config and problems files, so the fixed-vector-width kernels are the ones a full run has, and only what the list needs is built. A listed kernel is measured only if a full run would pair it with that problem; a `"*"` problem is measured exactly as a full run measures it, including its fallback to fixed-vector-width kernels when a native kernel fails to build. The files a run writes (`--csv`, `--failures-csv`, `--repeat-csv`) and the `--pairs` file must be distinct; the driver refuses to start otherwise. The README gains a "Measurement options" section. ## Unchanged by default Without the new options the driver behaves as before. Checked by running the develop driver and this one side by side on the same inputs, with stand-ins only for the compiler and the GPU: the CI config on rcr, rrr, ccr and crr, and `default_config.json` on rcr and crr, with problems that need fixed-vector-width kernels, some failed builds and some failed measurements. CSV bytes, the build list, every measured batch and the exit code were identical, and so was stdout apart from the elapsed benchmark time. The only difference reachable without the new options is the text of build-timeout messages. ## Testing - New CPU-only tests, registered in `dispatcher/tests/CMakeLists.txt`: `test_gemm_driver_measurement.py` (timeouts reach every step, failure records including build reasons and unsupported kernels, distinct output files, repeats, unchanged calls by default) and `test_gemm_driver_pairs.py` (`--list-pairs` equals the pairs a full run measures, `--pairs` builds and measures only the listed ones, `"*"` fallback, error cases). 43 tests pass. - Existing driver, MX and preshuffle tests give the same results as on develop: `test_gemm_driver_verify.py`, `test_gemm_driver_fp32.py`, `test_gemm_result_status.py`, `test_gemm_vector_sizes.py`, `test_dispatcher_cache_safety.py`, `test_mx_gemm_bridge.py`, `test_mx_gemm_cmake.py`, `test_preshuffle_bridge.py`, `test_preshuffle_gfx1250_pipelines.py`, `test_preshuffle_pipeline_matrix.py`, `test_quant_target_header_regressions.py` pass; `test_gemm_utils.py` has the same four failures on develop in this environment. - On `default_config.json` (rcr and crr), a `--list-pairs` file fed back as `--pairs` measures the same pairs and writes the same CSV as a full run; with injected build failures, all-`"*"` `--pairs` matches the full run pair for pair. - This version has not been run on a GPU: the change is in which pairs are built, launched, re-launched and recorded, not in the kernels or how they are timed. ## Merge order ```text this PR -> #13431 ``` #13431 carries the selector and the runbook, which use these options. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
illsilin
approved these changes
Oct 29, 2026
assistant-librarian Bot
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Oct 29, 2026
feat(dispatcher/heuristics): training data from the gemm driver's CSV (#13431) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit JIRA ID : AICK-2523 > This PR now also carries #13432 (kernel selection and the runbook), which was reviewed and approved there and merged into this branch. It has two commits: the training-data step, then #13432's. ## Why A full sweep builds every kernel of a configuration space and times it on every problem it can run, which can take days for a large space over many problems. A model trained on earlier measurements can rank each problem's candidates so that only its top K are built and timed, and a sweep of the same problems measures how good the picks are. This PR adds the two heuristics-side steps of that workflow, turning the tile_engine GEMM driver's results into training data and picking the kernels to measure, plus a runbook for the whole workflow ([HEURISTIC_BENCHMARKING.md](https://github.com/ROCm/rocm-libraries/blob/users/andriy/ck/heuristics-gemm-driver-csv/projects/composablekernel/dispatcher/heuristics/HEURISTIC_BENCHMARKING.md)). The driver-side options it uses are in #13430. Training data: the heuristic trains on canonical parquet rows (`data_pipeline.CANONICAL_COLUMNS`), and no in-tree tool converted the driver's CSV into them. `convert_csv_to_parquet.py`'s generic path matches only `gemm_universal_` names, so for the driver's `gemm_<dtype>_...` kernels it kept dtype and pipeline and lost tile, warps, warp tile, padding, epilogue and vector widths; it wrote non-canonical columns (`M`, `tflops`) and kept rows that failed `--verify` as labels. ## What changes Training data from the driver's CSV: - `data_pipeline.parse_gemm_driver_csv(path, arch)`: canonical rows from a `gemm_full_benchmark.py` CSV. Kernel fields, dtype and layout come from `parse_kernel_name` (vector widths included, 0 = native). Rows with `verified=False` are dropped and counted; an empty `verified` (no `--verify`) is kept. A name that does not decode raises. Floats are read round-trip. `bandwidth_gb_s` counts A and B in the input dtype and C in the driver's output dtype (fp16 for fp8/bf8). - `data_pipeline.py` accepts a `.csv` input. - `convert_csv_to_parquet.py` detects a driver CSV, delegates to it and adds the arch's `HW_PROFILES` constants. Batched GEMM results keep their own path, checked first; every other CSV keeps the generic path. Kernel selection and the runbook (from #13432): - `dispatcher/heuristics/select_kernels.py`: ranks the candidates from `gemm_full_benchmark.py --list-pairs` with `Predictor.rank_kernels` and writes the top K per problem as a `--pairs` file for the driver, plus a CSV of the picks (problem, rank, kernel, predicted TFLOP/s). - The candidates are the driver's own pairs, after its vector-width and tile rules, so every pick can launch. - Ties in the predicted value break on the kernel name. - The driver's field names map to the feature engine's (`wave_*` -> `warp_*`, `vector_size_*` -> `vec_*`). - It refuses a model trained for another dtype or arch, and non-universal variants. - `dispatcher/heuristics/HEURISTIC_BENCHMARKING.md`, a runbook: when the workflow applies, prerequisites and inputs, the reference sweep, training data and model, candidates and selection, measuring the picks, scoring them against the sweep, cost, and troubleshooting. It is linked from the heuristics README and the tile_engine GEMM README, and whitelisted in `dispatcher/heuristics/.gitignore` like the other tracked docs. - README: the conversion step under Quick Start. ## Testing - `tests/test_driver_csv.py` (10 tests): canonical columns, fields from names, the verify filter, exact floats, bandwidth for bf16 and fp8, hardware constants, the generic path for other CSVs, the `data_pipeline.py` CLI, and that `GemmUniversalVecFeatureEngine.extract_batch` reads the result. - `tests/test_select_kernels.py` (13 tests): - field mapping, top-K and tie-breaking, and the model/run checks; - an end-to-end run on the in-tree `gemm_universal_fp16_gfx950` model; - the runbook's scoring snippet, executed on constructed data; - every option the runbook's commands use exists in that script's argparse. The driver part skips until #13430 is in the tree. - Heuristics suite on develop with this PR: 576 passed, 2 skipped. - Real driver CSVs converted and trained with `train.py --operation gemm_universal_vec`: the model's spec records the vector-width engine, encoding version 1, 78 features and the arch's LDS capacity. - With #13430 merged in, the runbook's steps run as written. Training data, `train.py`, `--list-pairs` and `select_kernels.py` ran for real; the sweep and the picks run used stand-ins for the compiler and GPU. Only the picks' kernels were built, the measured pairs equalled the picks, and the scoring snippet ran on the outputs. ## Merge order ```text #13430 -> this PR ``` The runbook documents #13430's driver options, so #13430 merges first. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
Contributor
TheRock Submodule Bump ActivityNewest first
|
shumway
pushed a commit
to ROCm/composable_kernel
that referenced
this pull request
Nov 4, 2026
feat(dispatcher/heuristics): training data from the gemm driver's CSV (#13431) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit JIRA ID : AICK-2523 > This PR now also carries #13432 (kernel selection and the runbook), which was reviewed and approved there and merged into this branch. It has two commits: the training-data step, then #13432's. ## Why A full sweep builds every kernel of a configuration space and times it on every problem it can run, which can take days for a large space over many problems. A model trained on earlier measurements can rank each problem's candidates so that only its top K are built and timed, and a sweep of the same problems measures how good the picks are. This PR adds the two heuristics-side steps of that workflow, turning the tile_engine GEMM driver's results into training data and picking the kernels to measure, plus a runbook for the whole workflow ([HEURISTIC_BENCHMARKING.md](https://github.com/ROCm/rocm-libraries/blob/users/andriy/ck/heuristics-gemm-driver-csv/projects/composablekernel/dispatcher/heuristics/HEURISTIC_BENCHMARKING.md)). The driver-side options it uses are in #13430. Training data: the heuristic trains on canonical parquet rows (`data_pipeline.CANONICAL_COLUMNS`), and no in-tree tool converted the driver's CSV into them. `convert_csv_to_parquet.py`'s generic path matches only `gemm_universal_` names, so for the driver's `gemm_<dtype>_...` kernels it kept dtype and pipeline and lost tile, warps, warp tile, padding, epilogue and vector widths; it wrote non-canonical columns (`M`, `tflops`) and kept rows that failed `--verify` as labels. ## What changes Training data from the driver's CSV: - `data_pipeline.parse_gemm_driver_csv(path, arch)`: canonical rows from a `gemm_full_benchmark.py` CSV. Kernel fields, dtype and layout come from `parse_kernel_name` (vector widths included, 0 = native). Rows with `verified=False` are dropped and counted; an empty `verified` (no `--verify`) is kept. A name that does not decode raises. Floats are read round-trip. `bandwidth_gb_s` counts A and B in the input dtype and C in the driver's output dtype (fp16 for fp8/bf8). - `data_pipeline.py` accepts a `.csv` input. - `convert_csv_to_parquet.py` detects a driver CSV, delegates to it and adds the arch's `HW_PROFILES` constants. Batched GEMM results keep their own path, checked first; every other CSV keeps the generic path. Kernel selection and the runbook (from #13432): - `dispatcher/heuristics/select_kernels.py`: ranks the candidates from `gemm_full_benchmark.py --list-pairs` with `Predictor.rank_kernels` and writes the top K per problem as a `--pairs` file for the driver, plus a CSV of the picks (problem, rank, kernel, predicted TFLOP/s). - The candidates are the driver's own pairs, after its vector-width and tile rules, so every pick can launch. - Ties in the predicted value break on the kernel name. - The driver's field names map to the feature engine's (`wave_*` -> `warp_*`, `vector_size_*` -> `vec_*`). - It refuses a model trained for another dtype or arch, and non-universal variants. - `dispatcher/heuristics/HEURISTIC_BENCHMARKING.md`, a runbook: when the workflow applies, prerequisites and inputs, the reference sweep, training data and model, candidates and selection, measuring the picks, scoring them against the sweep, cost, and troubleshooting. It is linked from the heuristics README and the tile_engine GEMM README, and whitelisted in `dispatcher/heuristics/.gitignore` like the other tracked docs. - README: the conversion step under Quick Start. ## Testing - `tests/test_driver_csv.py` (10 tests): canonical columns, fields from names, the verify filter, exact floats, bandwidth for bf16 and fp8, hardware constants, the generic path for other CSVs, the `data_pipeline.py` CLI, and that `GemmUniversalVecFeatureEngine.extract_batch` reads the result. - `tests/test_select_kernels.py` (13 tests): - field mapping, top-K and tie-breaking, and the model/run checks; - an end-to-end run on the in-tree `gemm_universal_fp16_gfx950` model; - the runbook's scoring snippet, executed on constructed data; - every option the runbook's commands use exists in that script's argparse. The driver part skips until #13430 is in the tree. - Heuristics suite on develop with this PR: 576 passed, 2 skipped. - Real driver CSVs converted and trained with `train.py --operation gemm_universal_vec`: the model's spec records the vector-width engine, encoding version 1, 78 features and the arch's LDS capacity. - With #13430 merged in, the runbook's steps run as written. Training data, `train.py`, `--list-pairs` and `select_kernels.py` ran for real; the sweep and the picks run used stand-ins for the compiler and GPU. Only the picks' kernels were built, the measured pairs equalled the picks, and the scoring snippet ran on the outputs. ## Merge order ```text #13430 -> this PR ``` The runbook documents #13430's driver options, so #13430 merges first. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
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
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.
JIRA ID : AICK-2523
Why
A full sweep builds every kernel of a configuration space and times it on every problem it can run, which can take days for a large space over many problems. A model trained on earlier measurements can rank each problem's candidates so that only its top K are built and timed, and a sweep of the same problems measures how good the picks are. This PR adds the two heuristics-side steps of that workflow, turning the tile_engine GEMM driver's results into training data and picking the kernels to measure, plus a runbook for the whole workflow (HEURISTIC_BENCHMARKING.md). The driver-side options it uses are in #13430.
Training data: the heuristic trains on canonical parquet rows (
data_pipeline.CANONICAL_COLUMNS), and no in-tree tool converted the driver's CSV into them.convert_csv_to_parquet.py's generic path matches onlygemm_universal_names, so for the driver'sgemm_<dtype>_...kernels it kept dtype and pipeline and lost tile, warps, warp tile, padding, epilogue and vector widths; it wrote non-canonical columns (M,tflops) and kept rows that failed--verifyas labels.What changes
Training data from the driver's CSV:
data_pipeline.parse_gemm_driver_csv(path, arch): canonical rows from agemm_full_benchmark.pyCSV. Kernel fields, dtype and layout come fromparse_kernel_name(vector widths included, 0 = native). Rows withverified=Falseare dropped and counted; an emptyverified(no--verify) is kept. A name that does not decode raises. Floats are read round-trip.bandwidth_gb_scounts A and B in the input dtype and C in the driver's output dtype (fp16 for fp8/bf8).data_pipeline.pyaccepts a.csvinput.convert_csv_to_parquet.pydetects a driver CSV, delegates to it and adds the arch'sHW_PROFILESconstants. Batched GEMM results keep their own path, checked first; every other CSV keeps the generic path.Kernel selection and the runbook (from #13432):
dispatcher/heuristics/select_kernels.py: ranks the candidates fromgemm_full_benchmark.py --list-pairswithPredictor.rank_kernelsand writes the top K per problem as a--pairsfile for the driver, plus a CSV of the picks (problem, rank, kernel, predicted TFLOP/s).wave_*->warp_*,vector_size_*->vec_*).dispatcher/heuristics/HEURISTIC_BENCHMARKING.md, a runbook: when the workflow applies, prerequisites and inputs, the reference sweep, training data and model, candidates and selection, measuring the picks, scoring them against the sweep, cost, and troubleshooting. It is linked from the heuristics README and the tile_engine GEMM README, and whitelisted indispatcher/heuristics/.gitignorelike the other tracked docs.Testing
tests/test_driver_csv.py(10 tests): canonical columns, fields from names, the verify filter, exact floats, bandwidth for bf16 and fp8, hardware constants, the generic path for other CSVs, thedata_pipeline.pyCLI, and thatGemmUniversalVecFeatureEngine.extract_batchreads the result.tests/test_select_kernels.py(13 tests):gemm_universal_fp16_gfx950model;train.py --operation gemm_universal_vec: the model's spec records the vector-width engine, encoding version 1, 78 features and the arch's LDS capacity.train.py,--list-pairsandselect_kernels.pyran for real; the sweep and the picks run used stand-ins for the compiler and GPU. Only the picks' kernels were built, the measured pairs equalled the picks, and the scoring snippet ran on the outputs.Merge order
The runbook documents #13430's driver options, so #13430 merges first.
🤖 Generated with Claude Code