MLIR Release Notes
This document tries to provide some context about MLIR important changes in the context of LLVM releases. It is updated on a best effort basis.
At the moment the MLIR community does not qualify the LLVM release branch specifically, it is a snapshot of the MLIR development at the time of the release.
LLVM 24 ¶
Strict Property Assembly Formats ¶
Property-backed operations now use strict declarative assembly formats by
default. Setting the dialect option useStrictPropertiesInAssemblyFormat = 0
is deprecated and temporarily retains the legacy behavior. Dialects using
this setting must migrate their declarative assembly formats to bind every
inherent attribute and property directly or include prop-dict. Under the
default rules, attr-dict contains only discardable attributes. See
the guide.
GPU/AMDGPU Changes ¶
mlir::amdgpu::Chipsetis deprecated in favour ofmlir::ROCDL::TargetInfo, which describes a target by its triple, subarch, and the resolved set of target features from LLVM’s own tables. Lowerings should ask whether a target has a feature rather than inaccurately compare chipset versions.TargetInfoalso represents generic targets such asgfx9-4-genericand, unlikeChipset, explicitly stores the wavesize for targets where it is configurable.- The
chipsetoption in AMDGPU passes is renamed to anarchoption, which uses Clang target naming syntax. It accepts a GPU name with optional modifiers (gfx942,gfx942:xnack+,gfx9-4-generic), a triple (amdgpu9.42-amd-amdhsa), or a full target ID (amdgpu9.0a-amd-amdhsa--gfx90a:sramecc+:xnack-, which is whatrocminfoprints for a device’s ISA).chipsetorchipremain as compatibility names. The default arch isinvalid, so a target must be passed explicitly, removing the old “fallback”gfx000GPU. - Wavefront size is not a target-ID feature, so
convert-gpu-to-rocdltakes it as a separatewavesizeoption (32, 64, or 0 for the architecture’s default). Thewave64flag ongpu-lower-to-rocdl-pipelineand onrocdl-attach-targetis likewise replaced by the samewavesizeoption, which has the same allowed values. - In keeping with broader LLVM changes,
xnackandsrameccare no longer architecture features but module flags. In keeping with Clang, the target specifier still includes these xnack/sramecc flags where they’re configurable, but lowering passes now convert these to module flags. Downstream users should callmigrateArchFeaturesToModuleFlagsto lower these attributes in custom pipelines. rocdl-attach-targetgainsarchalongside its existingtriple,chipandfeatures. Whenarchis given, it overridestripleandchip, and handles xnack/sramecc modifier migration.gpu-lower-to-rocdl-pipeline, however, only supportsarch/chipnow, as manually setting a triple/chipset/features combination is a low-level operation that should require a manualrocdl-attach-targetcall.
Target-Independent GPU Transforms ¶
MLIRGPUTransforms no longer links any target dialect or target library. The
target-specific GPU passes and patterns moved into per-target libraries with
their own pass headers; pass names and options are unchanged:
| Library | Header | Contents |
|---|---|---|
MLIRGPUNVVMTransforms | mlir/Dialect/GPU/Transforms/NVVM/Passes.h | nvvm-attach-target |
MLIRGPUROCDLTransforms | mlir/Dialect/GPU/Transforms/ROCDL/Passes.h | rocdl-attach-target, populateGpuPromoteShuffleToAMDGPUPatterns, populateGpuLowerSubgroupReduceToDPPPatterns, populateGpuLowerClusteredSubgroupReduceToDPPPatterns |
MLIRGPUSPIRVTransforms | mlir/Dialect/GPU/Transforms/SPIRV/Passes.h | spirv-attach-target |
MLIRGPUXeVMTransforms | mlir/Dialect/GPU/Transforms/XeVM/Passes.h | xevm-attach-target |
Each header provides its own registration function (for example
registerGPUNVVMPasses()). The C API functions for these passes are renamed
to match, e.g. mlirCreateGPUGpuNVVMAttachTarget is now
mlirCreateGPUNVVMGpuNVVMAttachTarget.
LLVM 21 ¶
GPU/NVVM Changes ¶
- The default NVVM target architecture has been changed from
sm_50tosm_75.sm_75is the oldest GPU variant compatible with the widest range of recent major CUDA Toolkit versions (11/12/13). This affects theNVVMTargetAttr,GpuNVVMAttachTargetpass, and thegpu-lower-to-nvvm-pipeline.
LLVM 20 ¶
All the MLIR runners other than mlir-cpu-runner have been removed, as their functionality has been merged into it, and it has been renamed to mlir-runner.
LLVM 18 ¶
Properties: beyond attributes ¶
See LLVM 17 notes below. The Dialect option let usePropertiesForAttributes = 1; is
now the default. You can set it to 0 to revert to the previous behavior. This will be
removed in LLVM 19.
LLVM 17 ¶
See also the deprecations and refactoring doc.
Bytecode ¶
MLIR now support a bytecode serialization with versionning compatibility allowing 2 ways compatibility scheme, and lazy-loading capabilities.
Properties: beyond attributes ¶
This is a new mechanism to implement storage for operations without having to
use attributes. You can opt-in to use Properties for ODS inherent attributes
using let usePropertiesForAttributes = 1; in your dialect definition (the flag
will be default in the next release). See
slides and
recording of the open meeting presentation for
details.
Action: Tracing and Debugging MLIR-based Compilers ¶
Action is a new mechanism to encapsulate any transformation of any granularity in a way that can be intercepted by the framework for debugging or tracing purposes, including skipping a transformation programmatically (think about “compiler fuel” or “debug counters” in LLVM). As such, “executing a pass” is an Action, so is “try to apply one canonicalization pattern”, or “tile this loop”.
slides and recording of the open meeting presentation for details.
Transform Dialect ¶
See this EuroLLVM talk and the online tutorial.
Others ¶
- There is now support for “ distinct attributes”.
- “Resources” (a way to store data outside the MLIR context) and “configuration” can now be serialized alongside the IR.
MLIR