[mlir] (NFC) Clean up bazel and CMake target namesAll dialect targets in bazel have been named *Dialect and all dialecttargets in CMake have been named MLIR*Dialect.
[mlir] Refactor DialectRegistry delayed interface support into a general DialectExtension mechanismThe current dialect registry allows for attaching delayed interfaces, that are added to attrs/dial
[mlir] Refactor DialectRegistry delayed interface support into a general DialectExtension mechanismThe current dialect registry allows for attaching delayed interfaces, that are added to attrs/dialects/ops/etc.when the owning dialect gets loaded. This is clunky for quite a few reasons, e.g. each interface type has aseparate tracking structure, and is also quite limiting. This commit refactors this delayed mutation ofdialect constructs into a more general DialectExtension mechanism. This mechanism is essentially a registrationcallback that is invoked when a set of dialects have been loaded. This allows for attaching interfaces directlyon the loaded constructs, and also allows for loading new dependent dialects. The latter of which isextremely useful as it will now enable dependent dialects to only apply in the contexts in which theyare necessary. For example, a dialect dependency can now be conditional on if a user actually needs theinterface that relies on it.Differential Revision: https://reviews.llvm.org/D120367
show more ...
[mlir][nvvm] Fix bug in ldmatrix intrinsic conversionThe ldmatrix intrinsic trans option was inverted.Bug found by @christopherbate!Differential Revision: https://reviews.llvm.org/D121666
[mlir][NVVM] Add ldmatrix op to NVVM dialectDifferential Revision: https://reviews.llvm.org/D121347
Adjust "end namespace" comment in MLIR to match new agree'd coding styleSee D115115 and this mailing list discussion:https://lists.llvm.org/pipermail/llvm-dev/2021-December/154199.htmlDifferenti
Adjust "end namespace" comment in MLIR to match new agree'd coding styleSee D115115 and this mailing list discussion:https://lists.llvm.org/pipermail/llvm-dev/2021-December/154199.htmlDifferential Revision: https://reviews.llvm.org/D115309
[mlir][gpu] Extend shuffle op modes and add nvvm loweringAdd up, down and idx modes to gpu shuffle ops, also change the mode fromstring to enumDifferential Revision: https://reviews.llvm.org/D11
[mlir][gpu] Extend shuffle op modes and add nvvm loweringAdd up, down and idx modes to gpu shuffle ops, also change the mode fromstring to enumDifferential Revision: https://reviews.llvm.org/D114188
[mlir] Convert NamedAttribute to be a classNamedAttribute is currently represented as an std::pair, but thiscreates an extremely clunky .first/.second API. This commitconverts it to a class, with
[mlir] Convert NamedAttribute to be a classNamedAttribute is currently represented as an std::pair, but thiscreates an extremely clunky .first/.second API. This commitconverts it to a class, with better accessors (getName/getValue)and also opens the door for more convenient API in the future.Differential Revision: https://reviews.llvm.org/D113956
[mlir][nvvm] Generalize wmma ops to handle more types and shapeswmma intrinsics have a large number of combinations, ideally we want to be ableto target all the different variants. To avoid a comb
[mlir][nvvm] Generalize wmma ops to handle more types and shapeswmma intrinsics have a large number of combinations, ideally we want to be ableto target all the different variants. To avoid a combinatorial explosion in thenumber of mlir op we use attributes to represent the different variation ofload/store/mma ops. We also can generate with tablegen helpers to know whichcombinations are available. Using this we can avoid having too hardcode a pathfor specific shapes and can support more types.This patch also adds boiler plates for tf32 op support.Differential Revision: https://reviews.llvm.org/D112689
[MLIR][GPU][NVVM] Add warp synchronous matrix-multiply accumulate opsAdd warp synchronous matrix-multiply accumulate ops in GPU and NVVMdialect. Add following three ops to GPU dialect :- 1.) sub
[MLIR][GPU][NVVM] Add warp synchronous matrix-multiply accumulate opsAdd warp synchronous matrix-multiply accumulate ops in GPU and NVVMdialect. Add following three ops to GPU dialect :- 1.) subgroup_mma_load_matrix 2.) subgroup_mma_store_matrix 3.) subgroup_mma_computeAdd following three ops to NVVM dialect :- 1.) wmma.m16n16k16.load.[a,b,c].[f16,f32].row.stride 2.) wmma.m16n16k16.store.d.[f16,f32].row.stride 3.) wmma.m16n16k16.mma.row.row.[f16,f32].[f16,f32]Reviewed By: bondhugula, ftynse, ThomasRaouxDifferential Revision: https://reviews.llvm.org/D95330
[mlir] make implementations of translation to LLVM IR interfaces privateThere is no need for the interface implementations to be exposed, opaqueregistration functions are sufficient for all users,
[mlir] make implementations of translation to LLVM IR interfaces privateThere is no need for the interface implementations to be exposed, opaqueregistration functions are sufficient for all users, similarly to passes.Reviewed By: mehdi_aminiDifferential Revision: https://reviews.llvm.org/D97852
[mlir] add verifiers for NVVM and ROCDL kernel attributesMake sure they can only be attached to LLVM functions as a result of convertingGPU functions to the LLVM Dialect.
[mlir] Use the interface-based translation for LLVM "intrinsic" dialectsPort the translation of five dialects that define LLVM IR intrinsics(LLVMAVX512, LLVMArmNeon, LLVMArmSVE, NVVM, ROCDL) to th
[mlir] Use the interface-based translation for LLVM "intrinsic" dialectsPort the translation of five dialects that define LLVM IR intrinsics(LLVMAVX512, LLVMArmNeon, LLVMArmSVE, NVVM, ROCDL) to the new dialectinterface-based mechanism. This allows us to remove individual translationsthat were created for each of these dialects and just use one commonMLIR-to-LLVM-IR translation that potentially supports all dialects instead,based on what is registered and including any combination of translatabledialects. This removal was one of the main goals of the refactoring.To support the addition of GPU-related metadata, the translation interface isextended with the `amendOperation` function that allows the interfaceimplementation to post-process any translated operation with dialect attributesfrom the dialect for which the interface is implemented regardless of theoperation's dialect. This is currently applied to "kernel" functions, but canbe used to construct other metadata in dialect-specific ways withoutnecessarily affecting operations.Depends On D96591, D96504Reviewed By: nicolasvasilacheDifferential Revision: https://reviews.llvm.org/D96592