Commit Graph

10951 Commits

Author SHA1 Message Date
Alex Zinenko e5a5e00825 [mlir] fix compiler warnings
-Wsign-compare and -Wunsued-value in the recently introduced code.
2022-04-14 15:26:50 +02:00
Lei Zhang e54236dfb5 [mlir][vector] Cast away leading one dims for insert ops
Reviewed By: ThomasRaoux

Differential Revision: https://reviews.llvm.org/D123621
2022-04-14 08:57:32 -04:00
Lei Zhang bc408afbfe [mlir][vector] Fold splat constant transpose
Reviewed By: ThomasRaoux

Differential Revision: https://reviews.llvm.org/D123595
2022-04-14 08:51:25 -04:00
Alex Zinenko d064c4801c [mlir] Introduce Transform dialect
This dialect provides operations that can be used to control transformation of
the IR using a different portion of the IR. It refers to the IR being
transformed as payload IR, and to the IR guiding the transformation as
transform IR.

The main use case for this dialect is orchestrating fine-grain transformations
on individual operations or sets thereof. For example, it may involve finding
loop-like operations with specific properties (e.g., large size) in the payload
IR, applying loop tiling to those and only those operations, and then applying
loop unrolling to the inner loops produced by the previous transformations. As
such, it is not intended as a replacement for the pass infrastructure, nor for
the pattern rewriting infrastructure. In the most common case, the transform IR
will be processed and applied to payload IR by a pass. Transformations
expressed by the transform dialect may be implemented using the pattern
infrastructure or any other relevant MLIR component.

This dialect is designed to be extensible, that is, clients of this dialect are
allowed to inject additional operations into this dialect using the newly
introduced in this patch `TransformDialectExtension` mechanism. This allows the
dialect to avoid a dependency on the implementation of the transformation as
well as to avoid introducing dialect-specific transform dialects.

See https://discourse.llvm.org/t/rfc-interfaces-and-dialects-for-precise-ir-transformation-control/60927.

Reviewed By: nicolasvasilache, Mogball, rriddle

Differential Revision: https://reviews.llvm.org/D123135
2022-04-14 13:48:45 +02:00
Alex Zinenko 09141f1adf [mlir] Split intrinsics out of LLVMOps.td
Move the operations that correspond to LLVM IR intrinsics in a separate .td
file. This makes it easier to maintain the intrinsics and decreases the compile
time of LLVMDialect.cpp by ~25%.

Depends On D123310

Reviewed By: wsmoses, jacquesguan

Differential Revision: https://reviews.llvm.org/D123315
2022-04-14 13:23:31 +02:00
Alex Zinenko 2366a43b3c [mlir] initial support for opaque pointers in the LLVM dialect
LLVM IR has introduced and is moving forward with the concept of opaque
pointers, i.e. pointer types that are not carrying around the pointee type.
Instead, memory-related operations indicate the type of the data being accessed
through the opaque pointer. Introduce the initial support for opaque pointers
in the LLVM dialect:

  - `LLVMPointerType` to support omitting the element type;
  - alloca/load/store/gep to support opaque pointers in their operands and
    results; this requires alloca and gep to store the element type as an
    attribute;
  - memory-related intrinsics to support opaque pointers in their operands;
  - translation to LLVM IR for the ops above is no longer using methods
    deprecated in LLVM API due to the introduction of opaque pointers.

Unlike LLVM IR, MLIR can afford to support both opaque and non-opaque pointers
at the same time and simplify the transition. Translation to LLVM IR of MLIR
that involves opaque pointers requires the LLVMContext to be configured to
always use opaque pointers.

Reviewed By: wsmoses

Differential Revision: https://reviews.llvm.org/D123310
2022-04-14 13:23:29 +02:00
Mehdi Amini 35f48edb91 Apply clang-tidy fixes for llvm-qualified-auto in VectorTransforms.cpp (NFC) 2022-04-14 09:42:37 +00:00
Mehdi Amini 2d06420879 Apply clang-tidy fixes for performance-for-range-copy in SCF.cpp (NFC) 2022-04-14 09:42:36 +00:00
Uday Bondhugula 3766ca75f8 [MLIR] Fix missing return statement warning in PatternMatch.h
Fix missing return statement warning in PatternMatch.h. NFC.

```
mlir/include/mlir/IR/PatternMatch.h:983:3: warning: no return statement in
function returning non-void [-Wreturn-type]
```

Differential Revision: https://reviews.llvm.org/D123756
2022-04-14 13:03:21 +05:30
Arnab Dutta 392d55c1e2 [MLIR][GPU] Add canonicalization patterns for folding simple gpu.wait ops.
* Fold away redundant %t = gpu.wait async + gpu.wait [%t] pairs.

* Fold away %t = gpu.wait async ... ops when %t has no uses.

* Fold away gpu.wait [] ops.

* In case of %t1 = gpu.wait async [%t0], replace all uses of %t1
  with %t0.

Differential Revision: https://reviews.llvm.org/D121878
2022-04-14 12:30:55 +05:30
Stella Laurenzo 61352a580a [mlir] Introduce ml_program dialect.
Differential Revision: https://reviews.llvm.org/D120203
2022-04-13 21:38:14 -07:00
Stella Laurenzo 4f7585195d [NFC] Generically resolve body in FunctionOpInterface verifyBody.
Since the actual implementation class can arbitrarily shadow parts of the FunctionOpInterface exported API, access the body generically instead of via the use of the interface being defined.

Fixes https://github.com/llvm/llvm-project/issues/54807

Differential Revision: https://reviews.llvm.org/D123757
2022-04-13 20:58:19 -07:00
Arjun P 094ad0667c [MLIR][Presburger] change some `push_back`s to `emplace_back`s 2022-04-14 04:27:21 +01:00
Arjun P ef95a6e827 [MLIR][Presburger] change some post-increments/decrements to pre-increments/decrements 2022-04-14 04:26:29 +01:00
Okwan Kwon 18a01527ea [mlir] Fix a typo to load lsp-mode correctly.
Differential Revision: https://reviews.llvm.org/D123745
2022-04-14 00:04:40 +00:00
Mehdi Amini 93b23c8439 Apply clang-tidy fixes for llvm-else-after-return in SCF.cpp (NFC) 2022-04-13 23:14:44 +00:00
Mehdi Amini f19e90bc1f Apply clang-tidy fixes for readability-identifier-naming in OpenMPDialect.cpp (NFC) 2022-04-13 23:14:44 +00:00
Christopher Bate 77d2c815f5 [MLIR][GPU] Add GPU ops nvvm.mma.sync, nvvm.mma.ldmatrix, lane_id
This change adds three new operations to the GPU dialect: gpu.mma.sync,
gpu.mma.ldmatrix, and gpu.lane_id. The former two are meant to target
the lower level nvvm.mma.sync and nvvm.ldmatrix instructions, respectively.
Lowerings are added for the new GPU operations for conversion to
NVVM.

Reviewed By: ThomasRaoux

Differential Revision: https://reviews.llvm.org/D123647
2022-04-13 22:50:07 +00:00
Thomas Raoux 5b1b7108c8 [mlir][vector] Add unrolling pattern for TransposeOp
Support unrolling for vector.transpose following the same interface as
other vector unrolling ops.

Differential Revision: https://reviews.llvm.org/D123688
2022-04-13 19:44:16 +00:00
Mogball 60e34f8ddd [mlir][ods] Remove StrEnumAttr
StrEnumAttr has been deprecated in favour of EnumAttr, a solution based on AttrDef (https://reviews.llvm.org/D115181). This patch removes StrEnumAttr, along with all the custom ODS logic required to handle it.

See https://discourse.llvm.org/t/psa-stop-using-strenumattr-do-use-enumattr/5710 on how to transition to EnumAttr. In short,

```
// Before
def MyEnumAttr : StrEnumAttr<"MyEnum", "", [
  StrEnumAttrCase<"A">,
  StrEnumAttrCase<"B">
]>;

// After (pick an integer enum of your choice)
def MyEnum : I32EnumAttr<"MyEnum", "", [
  I32EnumAttrCase<"A", 0>,
  I32EnumAttrCase<"B", 1>
]> {
  // Don't generate a C++ class! We want to use the AttrDef
  let genSpecializedAttr = 0;
}
// Define the AttrDef
def MyEnum : EnumAttr<MyDialect, MyEnum, "my_enum">;
```

Reviewed By: rriddle, jpienaar

Differential Revision: https://reviews.llvm.org/D120834
2022-04-13 17:49:02 +00:00
Jacques Pienaar 5a236e69f2 [mlir][pdll] Include string in PDLLServer.h 2022-04-13 10:31:06 -07:00
Jacques Pienaar a9d7598ef8 [mlir][pdll] Rename extra dir flag
Not sure why this is overlapping, need to repro locally, but just rename
to something more explicit instead.
2022-04-13 10:00:10 -07:00
Jacques Pienaar aa9da11f5d [mlir][pdll] Add extra-dirs for LSP includes.
Enable specifying additional include directories to search. This is
consistent with what one can do with clangd (although there it is more
general compilation options) and Python LSP. We would in general expect
these to be provided by compilation database equivalent.

Differential Revision: https://reviews.llvm.org/D123474
2022-04-13 09:41:45 -07:00
Marius Brehler f71f105901 [mlir][docs] Fix broken links 2022-04-13 16:18:20 +02:00
Adrian Kuegel bd941c83f5 [mlir][Tensor] Fix wrong comment (NFC) 2022-04-13 14:30:55 +02:00
Shraiysh Vaishay f0ba32d666 [mlir][LLVM-IR] Added support for global variable attributes
This patch adds thread_local to llvm.mlir.global and adds translation for dso_local and addr_space to and from LLVM IR.

Reviewed By: Mogball

Differential Revision: https://reviews.llvm.org/D123412
2022-04-13 08:31:01 +05:30
jacquesguan 605fc89a61 [mlir][Arithmetic] Add common constant folder function for type cast ops.
This revision replaces current type cast constant folder with a new common type cast constant folder function template.
It will cover all former folder and support fold the constant splat and vector.

Differential Revision: https://reviews.llvm.org/D123489
2022-04-13 02:11:59 +00:00
Aart Bik 28063a281b [mlir][sparse] refactored python setup of sparse compiler
Reviewed By: bixia

Differential Revision: https://reviews.llvm.org/D123419
2022-04-12 11:58:41 -07:00
Mahesh Ravishankar b40e901333 [mlir][Linalg] Allow collapsing subset of the reassociations when fusing by collapsing.
This change generalizes the fusion of `tensor.expand_shape` ->
`linalg.generic` op by collapsing to handle cases where only a subset
of the reassociations specified in the `tensor.expand_shape` are valid
to be collapsed.
The method that does the collapsing is refactored to allow it to be a
generic utility when required.

Reviewed By: gysit

Differential Revision: https://reviews.llvm.org/D123153
2022-04-12 18:56:32 +00:00
Shraiysh Vaishay b18e82186f [mlir][OpenMP] Added omp.task
This patch adds tasking construct according to Section 2.10.1 of OpenMP 5.0

Reviewed By: peixin, kiranchandramohan, abidmalikwaterloo

Differential Revision: https://reviews.llvm.org/D123575
2022-04-12 23:55:47 +05:30
Groverkss 20aedb148b [MLIR][Presburger] Remove inheritance from PresburgerSpace in IntegerRelation, PresburgerRelation and PWMAFunction
This patch removes inheritence from PresburgerSpace in IntegerRelation and
instead makes it a member of these classes.

This is required for three reasons:
  - It prevents implicit casting to PresburgerSpace.
  - Not all functions of PresburgerSpace need to be exposed by the deriving classes.
  - IntegerRelation and IntegerPolyhedron are defined in a PresburgerSpace. It
    makes more sense for the space to be a member instead of them inheriting from
    a space.

Reviewed By: arjunp, ftynse

Differential Revision: https://reviews.llvm.org/D123585
2022-04-12 22:48:52 +05:30
Arjun P 0ac213667d [MLIR][Presburger] normalizeDiv: add assert that denom > 0 2022-04-12 13:06:53 +01:00
jacquesguan 83bd4fe2e8 [mlir][Math] Replace some constant folder functions with common folder functions.
Differential Revision: https://reviews.llvm.org/D123485
2022-04-12 11:34:47 +00:00
Arjun P 4aeb2a57f4 [MLIR][Presburger][Simplex] addSymbolicCut: don't add symbol div if denom is 1
This is unncessary, so we remove it as an optimization.

Reviewed By: Groverkss

Differential Revision: https://reviews.llvm.org/D123540
2022-04-12 12:27:27 +01:00
Andrzej Warzynski fb16ed258c [mlir] Prefix pass manager options with `mlir-`
With this change, there's going to be a clear distinction between LLVM
and MLIR pass maanger options (e.g. `-mlir-print-after-all` vs
`-print-after-all`). This change is desirable from the point of view of
projects that depend on both LLVM and MLIR, e.g. Flang.

For consistency, all pass manager options in MLIR are prefixed with
`mlir-`, even options that don't have equivalents in LLVM .

Differential Revision: https://reviews.llvm.org/D123495
2022-04-12 09:32:44 +00:00
Matthias Springer fa087b4352 [mlir][scf][bufferize][NFC] Lookup buffer using helper function
Lookup iter_arg buffers using `lookupBuffer` instead of always creating a new `ToMemrefOp`. Also cast all yielded buffers (if necessary), regardless of whether they are an equivalent buffer or a new allocation.

Note: This should have been part of D123369.

Differential Revision: https://reviews.llvm.org/D123383
2022-04-12 18:09:30 +09:00
Mehdi Amini eeef40a6cd Apply clang-tidy fixes for performance-unnecessary-value-param in LLVMDialect.cpp (NFC) 2022-04-12 07:47:04 +00:00
Mehdi Amini e188ad8bf0 Apply clang-tidy fixes for performance-unnecessary-value-param in SplitReduction.cpp (NFC) 2022-04-12 07:47:04 +00:00
Mehdi Amini 30846d2916 Guard copy of std::function to llvm::function_ref (fix crash)
This is a footgun: assigning a null std::function to a function_ref
does not yield a null function_ref...
2022-04-12 07:44:04 +00:00
Mehdi Amini 07db69effe Use std::function instead of function_ref in MLIR JitRunner
This fixes an ASAN failure.
2022-04-12 07:28:19 +00:00
Mehdi Amini 6b7e6ea489 Revert "Fix CUDA runtime wrapper for GPU mem alloc/free to async"
This reverts commit b4117fede2.
This broke one of the MLIR bot, a test is failing.
2022-04-12 06:50:27 +00:00
Mehdi Amini 3c9e09036c Apply clang-tidy fixes for readability-identifier-naming in LinalgOps.cpp (NFC) 2022-04-12 05:30:01 +00:00
Mehdi Amini 47286e17ed Apply clang-tidy fixes for performance-for-range-copy in LinalgOps.cpp (NFC) 2022-04-12 05:30:01 +00:00
Mehdi Amini 0149a437e2 Apply clang-tidy fixes for llvm-qualified-auto in LinalgOps.cpp (NFC) 2022-04-12 04:39:58 +00:00
Mehdi Amini bf62a4b9c5 Apply clang-tidy fixes for performance-move-const-arg in ArithmeticOps.cpp (NFC) 2022-04-12 04:39:57 +00:00
Uday Bondhugula 56245cc18c [MLIR] NFC. Address clang-tidy warning in AffineOps.cpp
NFC. Address clang-tidy warning in AffineOps.cpp.
2022-04-12 09:58:29 +05:30
Uday Bondhugula b4117fede2 Fix CUDA runtime wrapper for GPU mem alloc/free to async
Switch CUDA runtime wrapper for GPU mem alloc/free to async. The
semantics of the GPU dialect ops (gpu.alloc/dealloc) and the wrappers it
lowered to (gpu-to-llvm) was for the async versions -- however, this was
being incorrectly mapped to cuMemAlloc/cuMemFree instead of
cuMemAllocAsync/cuMemFreeAsync.

Reviewed By: csigg

Differential Revision: https://reviews.llvm.org/D123482
2022-04-12 09:04:02 +05:30
Peixin-Qiao 139c5b8a96 [MLIR][OpenMP] Add support for threadprivate directive
This supports the threadprivate directive in OpenMP dialect following
the OpenMP 5.1 [2.21.2] standard. Also lowering to LLVM IR using OpenMP
IRBduiler.

Reviewed By: kiranchandramohan, shraiysh, arnamoy10

Differential Revision: https://reviews.llvm.org/D123350
2022-04-12 10:15:15 +08:00
jacquesguan 920527f4b5 [mlir][NFC] Remove some redundant code.
Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D123487
2022-04-12 02:10:04 +00:00
Eugene Zhulenev b35b9e307f [mlir] Add msan memory unpoisoning macros to mlir ExecutionEngine
Adding annotations on as-needed bases, currently only for memrefCopy, but in general all C API functions that take pointers to memory allocated/initialized inside the jit-compiled code must be annotated, to be able to run with msan.

Reviewed By: mehdi_amini

Differential Revision: https://reviews.llvm.org/D123557
2022-04-11 18:58:28 -07:00