Reduction fusion testing. - #5
Closed
charithaintc wants to merge 68 commits into
Closed
Conversation
Also fixed the math.erf issue with convert-math-to-libm after convert-math-to-llvm to make sure we only apply to those that can't be lowered directly to intrinsics. Ongoing efforts to llvm#142
This PR enables a new mode in kernel_bench that allows one to use the Torch Dynamo compiler plugin (CPU only for now). This is similar to the previous imported module: it runs the same pipeline, compiles the same MLIR modules and yield the same results. Major changes: * Factoring some code around import & pipeline - import_module extracted in lh.ingress.torch - make_function_callable an external helper now * Make kernel_bench import the Torch model first - Run torch eager to gather the reference output - import to MLIR or run with torch.compile - reuse the same pipeline on both cases - Compare outputs instead of printing output * Allowing empty input/output shapes for torch-compile - For Dynamo this can be fetched from model and used - Import can probably also, but for a later PR Torch Compile mode enabled in CI with the previous tests. Progress on llvm#142 Note: This is not the default yet due to benchmark not enabled on this mode yet. A follow up PR will add support and enable it by default.
Gets rid of the `memrefCopy` error on some kernels. See llvm#142
Fixed a `triton.language` mismatch.
* Cleanup on unused args in torch_compile * New flag `validate` to skip eager run on IR investigations * Passing unknown test_kernel_bench args to kernel_bench * Rename --test to --kernel * Remove random seed to increase coverage in CI * Warning when --benchmark is selected but no flops given
There was a mismatch between the number of argument FX was expecting and the arguments the MLIR code was returning, added by the MLIR backend. This PR corrects that imbalance and returns the correct number of arguments. Tests adapted to avoid OOM-kill of some very large kernels. The allocation of additional multi-GB tensors for return values is not reasonable, we need to change that approach in a future PR. Fixes the last runtime problem in llvm#142
Allow user to bypass target detection and force arch/feature. If there's a schedule that matches that arch/feature and data type, it will use to call kernel_bench with it. Otherwise, bail. If both arch and feature provided, try feature first, then arch, then bail to default schedule. If only feature is provided, auto-detect arch and try that if feature schedule doesn't exist. Fixes llvm#141
This will give us 2 months more before they remove the old builds.
Some sizes were too large for CPU execution, especially when importing the models to file allocating all inputs, but also when using torch.compile. This PR reduces many test sizes to run on the smoke test. The original sizes can stil be used by omitting the --input-shapes and --output-shape arguments for torch.compile. Tested with both import and torch-compile. This has the minimum set of failures we have seen so far.
This PR allows one to pass the Module init arguments through the command line, so that we can match the sizes of the execution with the internal sizes of the module (as we change them). Key changes: * Allow kernel_bench to override the init_args * Allow test_kernel_bench to read from yaml and pass to kernel_bench Minor change: * Factor out type parser into helper, allow tuples
Extends the XeGPU MLP/matmul schedule with a cost model that can be used to generate valid tile sizes for given (M, N, K) matmul shape.
…m#168) Also, rename `kernel_bench` to `kernel-bench` and `test_kernel_bench` to `test-kernel-bench` to match the other tools (and existing documentation).
Removing old examples that got replaced with better ones. Renaming directories to make it easier to understand the examples. XeGPU examples untouched, for a follow up PR.
Allows the Runner to output compiled object file. It can be further processed with `objdump` to produce readable assembly. Fixes llvm#164
Doesn't seem to be used anywhere, now that the Runner has taken over.
…file (llvm#169) Adds an optional `model_datatype` argument to torch `import_model` and `import_from_file` methods. Setting `model_datatype=torch.float16` converts all the hidden model weights etc. to float16.
Needed after directory move.
`tune_matmul_ga.py` now takes gpu target as an argument and uses `XeGPUSpecs`.
As noted in the TODO, this PR makes use of llvm#169 to allow conversion of models to an appropriate data type. For now, only F32 and BF16 are supported, but it should not be hard to make other data types available. Also, fix the bf16 output return and add tests in CI.
Reaching twice the performance of compiler vectorized code, and about 2/3 of peak for matmul, so not bad, but also nothing special. This is basically a single change from the matmul ones, so will serve as a way to test parametrization next more than anything, but already makes the CI faster.
After llvm#174 we can now dump the object file, so we use that to disassemble and dump the assembly to a file with the same name as the input, but with a .s extension. Also simplify some repeated `if benchmark` logic in the Runner.
Refactors PyTorch MLIR backend to lighthouse's Runner for consistent execution path. Assisted-by: Claude
Adjusts the default YAML pipeline file name.
Reduce some noise from benchmarks.
update status with a more visual approach
Enables MLIR object file creation in the PyTorch backend. Removes temporary file creation for the object files. Instead, users must explicitly provide a file name for the dumps. Tools are updated accordingly. The new functionality is used in kernel-bench tool to enable assembly dumps when torch-compile option is enabled.
This PR adds support for variables in descriptor files. It uses the same mechanism for passing options to schedules and the existing syntax so the only difference is to add a variable replacement in PipelineDescriptor's constructor to parse and update. This simplifies considerably the pipelines in the kernel bench testing framework and allows for choices to be propagated through the includes to build trully parametric pipelines from one set of top-level arguments.
- Updates mlp schedule to support transpose a and/or transpose b case. - `inspect_payload` returns metadata in generic `"layers"` nested dict. Currently returns metadata for matmul, batch_matmul and elemwise layers. - Matmul cost model and parameter selector are updated to handle transpose cases. - Cost model `generate_configs` has a verbose flag which prints tile selection info. - `get_tileable_consumers` returns only matmul epilog, i.e., excludes next linalg.matmul op. - matmul and mlp examples are updated to have `--transpose-a/b` options.
Improves python transforms importer to handle relative imports present in python files defining MLIR transforms. The file importers are moved into a standalone utility submodule. Assisted-by: Claude
Updates tools usage examples with newest changes to their arguments.
Follow up from llvm#193, now we can run benchmarks inside torch-compile mode
Updates PyTorch to 2.12.1 release. Since 2.10 PyTorch packages Triton under its default 'triton' name. Thus, the explicit 'pytorch_triton' dependency is no longer needed.
) - Updates MLP schedule constraints: SG tile size can be 16 (minimum was 32) + bugfixes. - `tune_costmodel.py` is refactored to have a more generic `optimize_payload` entry point. - Removes `tune_matmul_gridsearch.py`, `tune_matmul_ga.py`, `genetic_algorithm.py`. - Updates llvm version to the latest one. This includes upstream xegpu support for transpose-A gemms etc.
…vm#199) In the matmul and MLP examples, A and B datatype can be set to bfloat16 instead of default float16. Lowering and execution works without any additional changes.
BF16 conversion was a bit broken and not correctly propagated to the `import_torch` function. Co-authored-by: Adam Siemieniuk <adam.siemieniuk@intel.com>
- Adds a simple linalg-to-binary schedule for elementwise kernels. - Refactors common parts with the mlp schedule to sub-schedules. - The mlp and elementwise schedules now take the payload function name for more robust matching.
Adds a separate kernel config file for CI mode runs. The CI test spec run fewer kernels and uses smaller shapes for faster test completion.
… ops (llvm#205) Adds a utility transform op that filters out Linalg operations that have fewer number of loops than specified. The transform supplements tiling schedule which normally fails when too many tile sizes are provided. This way the tiling schedule can be applied more aggressively without halting lowering due to application failures when an op cannot be tiled anyway.
Adds examples/xegpu/kernel_bench.py routine to execute KernelBench benchmarks on Intel GPUs.
Adds a support transform and schedule that moves base memref offsets into a separate subview to align IR shape to AMX pattern expectations. Co-authored-by: Julian Oppermann <julian.oppermann@intel.com>
Simplify xegpu kernel bench CI tests.
Adds utility function that enables AMX registers in OS. Explicit AMX register config is required before AMX instructions can be executed.
- Adds a simplified xegpu matmul schedule that takes tile size parameters as kwargs. This is easier to hook up with the yaml desciptors. - Adds xegpu target support to `kernel-bench` tool. Currently enabled only in the torch.compile code path. - `test_kernel_bench.py` example: - Add `--pipeline` arg to optionally override the schedule yaml file look up. - Adds `schedules/xegpu/matmul/bf16.yaml` example schedule for a ~4k matmul. Compatible with the existing schedule lookup scheme.
Extends KernelBench runner to prioritize target and feature specific pipelines when available. The existing KB schedules are reorganized into common f32 and AMX-specific bf16 pipelines. KB shapes are updated to be more representative across currently supported extensions. --------- Co-authored-by: Renato Golin <rengolin@systemcall.eu>
Suppresses future warning coming from torch-mlir during model import.
Adds helper to find OpenMP library and a new schedule that injects OpenMP calls over parallel loops. The new parallelization schedule is integrated into x86 KernelBench pipelines to enable multicore execution. Assisted-by: Claude
Now that benchmarking is working for torch-compile, and that it's more stable than full conversion (less fiddling with arguments), we should have it as our default mode.
Batched script to execute a number of kernels and capture a lot of information on the execution for _every kernel_: * `kernel-bench` command line, as driven by `test-kernel-bench.py` * Full MLIR pipeline dump, into a log file * Object and assembly files * Validation against PyTorch, benchmark numbers, etc. The main point of this script is to help remote execution in a batched way, for example when SSH/SRUN on remote compute nodes, or multiple nodes, for wide collection of data and benchmark results. The secondary point is to help debug individual compilations by providing as much information as possible for various choices (different cmd-line args) and allowing one to compare them offline. Also, update kernel-bench's ability to receive names for the object and assembly files directly, to make it easier to automate and not pollute the repo or the KernelBench submodule.
Relaxes path normalization to first check if provided include path exists and then fallback to checking the default search locations. This allows including descriptors using relative paths e.g.: 'lh-opt --stage ./path/to/file.yaml'
charithaintc
marked this pull request as ready for review
July 22, 2026 19:38
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
5 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.
No description provided.