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.
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.
mlir::amdgpu::Chipset is deprecated in favour of mlir::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. TargetInfo also represents generic targets such as gfx9-4-generic and, unlike Chipset, explicitly stores the wavesize for targets where it is configurable.chipset option in AMDGPU passes is renamed to an arch option, 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 what rocminfo prints for a device's ISA). chipset or chip remain as compatibility names. The default arch is invalid, so a target must be passed explicitly, removing the old “fallback” gfx000 GPU.convert-gpu-to-rocdl takes it as a separate wavesize option (32, 64, or 0 for the architecture's default). The wave64 flag on gpu-lower-to-rocdl-pipeline and on rocdl-attach-target is likewise replaced by the same wavesize option, which has the same allowed values.xnack and sramecc are 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 call migrateArchFeaturesToModuleFlags to lower these attributes in custom pipelines.rocdl-attach-target gains arch alongside its existing triple, chip and features. When arch is given, it overrides triple and chip, and handles xnack/sramecc modifier migration.gpu-lower-to-rocdl-pipeline, however, only supports arch/chip now, as manually setting a triple/chipset/features combination is a low-level operation that should require a manual rocdl-attach-target call.sm_50 to sm_75. sm_75 is the oldest GPU variant compatible with the widest range of recent major CUDA Toolkit versions (11/12/13). This affects the NVVMTargetAttr, GpuNVVMAttachTarget pass, and the gpu-lower-to-nvvm-pipeline.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.
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.
See also the deprecations and refactoring doc.
MLIR now support a bytecode serialization with versionning compatibility allowing 2 ways compatibility scheme, and lazy-loading capabilities.
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 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.
See this EuroLLVM talk and the online tutorial.