Commit Graph

8164 Commits

Author SHA1 Message Date
Victor Perez
b8064c1a03 [mlir][llvm] Add trap intrinsics to the dialect
Define llvm.trap, llvm.debugtrap, and llvm.ubsantrap intrinsics in the
llvm dialect.

Signed-off-by: Victor Perez <victor.perez@codeplay.com>

Differential Revision: https://reviews.llvm.org/D149574
2023-05-02 10:09:31 +01:00
Mehdi Amini
5e118f933b Introduce MLIR Op Properties
This new features enabled to dedicate custom storage inline within operations.
This storage can be used as an alternative to attributes to store data that is
specific to an operation. Attribute can also be stored inside the properties
storage if desired, but any kind of data can be present as well. This offers
a way to store and mutate data without uniquing in the Context like Attribute.
See the OpPropertiesTest.cpp for an example where a struct with a
std::vector<> is attached to an operation and mutated in-place:

struct TestProperties {
  int a = -1;
  float b = -1.;
  std::vector<int64_t> array = {-33};
};

More complex scheme (including reference-counting) are also possible.

The only constraint to enable storing a C++ object as "properties" on an
operation is to implement three functions:

- convert from the candidate object to an Attribute
- convert from the Attribute to the candidate object
- hash the object

Optional the parsing and printing can also be customized with 2 extra
functions.

A new options is introduced to ODS to allow dialects to specify:

  let usePropertiesForAttributes = 1;

When set to true, the inherent attributes for all the ops in this dialect
will be using properties instead of being stored alongside discardable
attributes.
The TestDialect showcases this feature.

Another change is that we introduce new APIs on the Operation class
to access separately the inherent attributes from the discardable ones.
We envision deprecating and removing the `getAttr()`, `getAttrsDictionary()`,
and other similar method which don't make the distinction explicit, leading
to an entirely separate namespace for discardable attributes.

Recommit d572cd1b06 after fixing python bindings build.

Differential Revision: https://reviews.llvm.org/D141742
2023-05-01 23:16:34 -07:00
Aart Bik
9a018a7b48 [mlir][sparse] relax constraints on tensor.cast with pre-rewriting
Reviewed By: wrengr

Differential Revision: https://reviews.llvm.org/D149489
2023-05-01 16:03:44 -07:00
Mehdi Amini
1e853421a4 Revert "Introduce MLIR Op Properties"
This reverts commit d572cd1b06.

Some bots are broken and investigation is needed before relanding.
2023-05-01 15:55:58 -07:00
Mehdi Amini
d572cd1b06 Introduce MLIR Op Properties
This new features enabled to dedicate custom storage inline within operations.
This storage can be used as an alternative to attributes to store data that is
specific to an operation. Attribute can also be stored inside the properties
storage if desired, but any kind of data can be present as well. This offers
a way to store and mutate data without uniquing in the Context like Attribute.
See the OpPropertiesTest.cpp for an example where a struct with a
std::vector<> is attached to an operation and mutated in-place:

struct TestProperties {
  int a = -1;
  float b = -1.;
  std::vector<int64_t> array = {-33};
};

More complex scheme (including reference-counting) are also possible.

The only constraint to enable storing a C++ object as "properties" on an
operation is to implement three functions:

- convert from the candidate object to an Attribute
- convert from the Attribute to the candidate object
- hash the object

Optional the parsing and printing can also be customized with 2 extra
functions.

A new options is introduced to ODS to allow dialects to specify:

  let usePropertiesForAttributes = 1;

When set to true, the inherent attributes for all the ops in this dialect
will be using properties instead of being stored alongside discardable
attributes.
The TestDialect showcases this feature.

Another change is that we introduce new APIs on the Operation class
to access separately the inherent attributes from the discardable ones.
We envision deprecating and removing the `getAttr()`, `getAttrsDictionary()`,
and other similar method which don't make the distinction explicit, leading
to an entirely separate namespace for discardable attributes.

Differential Revision: https://reviews.llvm.org/D141742
2023-05-01 15:35:48 -07:00
Peiming Liu
d4db528938 [mlir][sparse] extend unpack operation to support unpacking a batched COO type
Reviewed By: aartbik

Differential Revision: https://reviews.llvm.org/D149103
2023-05-01 18:17:29 +00:00
Jacques Pienaar
5c90e1ffb0 [mlir][bytecode] Return error instead of min version
Can't return a well-formed IR output while enabling version to be bumped
up during emission. Previously it would return min version but
potentially invalid IR which was confusing, instead make it return
error and abort immediately instead.

Differential Revision: https://reviews.llvm.org/D149569
2023-04-30 22:11:02 -07:00
max
0a02f76d11 [MLIR][tensor] generate default builder for FromElementsOp
Removed builder is the same as default builder, with the added benefit that python bindings will be generated for the default builder.

Reviewed By: ftynse

Differential Revision: https://reviews.llvm.org/D149508
2023-04-30 18:44:52 -05:00
Jakub Kuderski
ef64e5913b [mlir][vector][NFC] Update vector.flat_transpose description
Change the example to match how this op is actually printed. Do not use
'smart quotes' in the description.

Reviewed By: jpienaar

Differential Revision: https://reviews.llvm.org/D149535
2023-04-29 18:34:34 -04:00
Jacques Pienaar
0610e2f6a2 [mlir][bytecode] Allow client to specify a desired version.
Add method to set a desired bytecode file format to generate. Change
write method to be able to return status including the minimum bytecode
version needed by reader. This enables generating an older version of
the bytecode (not dialect ops, attributes or types). But this does not
guarantee that an older version can always be generated, e.g., if a
dialect uses a new encoding only available at later bytecode version.
This clamps setting to at most current version.

Differential Revision: https://reviews.llvm.org/D146555
2023-04-29 05:35:53 -07:00
Jakub Kuderski
09043a26c8 [mlir][arith] Add patterns to commute extension over vector extraction
This moves zero/sign-extension ops closer to their use and exposes more
narrowing optimization opportunities.

Reviewed By: antiagainst

Differential Revision: https://reviews.llvm.org/D149233
2023-04-28 13:48:50 -04:00
Matthias Springer
77124386fe [mlir][tensor] Add transform to make tensor.pad loop-independent
Add a transform to make `tensor.pad` and `tensor.empty` ops independent of SCF loop IVs. Such ops can then be hoisted.

E.g.:
```
scf.for %iv = %lb to %ub step %step {
  %high = affine.apply affine_map<(d0)[s0] -> (s0 - d0)> (%i)[%ub]
  %p = tensor.pad %t low[5] high[%high] ...
  ...
}
```
Is transformed to:
```
%high_new = affine.apply affine_map<()[s0, s1] -> (-s0 + s1)> ()[%lb, %ub]
%p_hoistable = tensor.pad %t low[5] high[%high_new]
%dim = tensor.dim %t, %c0
%size = affine.apply affine_map<(d0)[s0, s1] -> (-d0 + s0 + s1 + 5)>(%iv)[%ub, %dim]
%slice = tensor.extract_slice %p_hoistable [0] [%size] [1]
```

Differential Revision: https://reviews.llvm.org/D143910
2023-04-28 11:46:32 +09:00
Andrew Gozillon
f478721231 [MLIR][LLVM] Add accessor for LLVMModule and invoke convertDialectAttributes on GlobalOps
This patch seeks to do two things add an accessor method to
retrieve the ModuleTranslations contained LLVM Module for direct
usage by dialects that are being lowered to LLVM-IR. One particular
use case for this is in the OpenMP Dialect, when interfacing
with the OMPIRBuilder in certain cases it is useful to be able
to access the LLVM Module directly.

The second is invoking convertDialectAttributes on GlobalOp's
so as to be able to lower dialect specific attributes that are
applied or lowered onto GlobalOp's.

Reviewers: ftynse

Differential Revision: https://reviews.llvm.org/D149279
2023-04-27 14:51:54 -05:00
Kiung Jung
eb22971bfb [MLIR][EmitC] Add empty emitc.constant check
Implementing logic to check if the emitc dialect constant Op is empty.

Reviewed By: marbre

Differential Revision: https://reviews.llvm.org/D147907
2023-04-27 16:03:53 +00:00
Kiung Jung
2aa1818a76 Disallow to apply the op & to const
Disallow to apply the operator & (address of) to emitc.constant operations.

Reviewed By: marbre, simon-camp

Differential Revision: https://reviews.llvm.org/D147992
2023-04-27 15:44:16 +00:00
gilsaia
e63cc56d9d [MLIR][presburger] normalize divisionrepr
Added a simple normalize function to divisionrepr and added a simple unittest.
Added a normalizediv call to divisionrepr's removeDuplicateDivs function, which now eliminates divs that are consistent after gcd's normalize

Reviewed By: Groverkss

Differential Revision: https://reviews.llvm.org/D147381
2023-04-27 14:13:15 +05:30
Kazu Hirata
1b2247d932 [mlir] Replace None with std::nullopt in comments (NFC)
This is part of an effort to migrate from llvm::Optional to
std::optional:

https://discourse.llvm.org/t/deprecating-llvm-optional-x-hasvalue-getvalue-getvalueor/63716
2023-04-26 23:29:02 -07:00
Théo Degioanni
f88f8fd0bc [mlir] Add a generic mem2reg implementation.
This patch introduces a generic implementation of mem2reg on
unstructured control-flow, along with a specialization for LLVM IR. This
is achieved by defining three new interfaces, representing 1. allocating
operations, 2. operations doing memory accesses, 3. operations that can
be rewired and/or deleted to stop using a specific use.

The file containing the core implementation of the algorithm
(`Mem2Reg.cpp`) contains a detailed explanation of how the algorithm
works. The contract for this pass is that given a memory slot with a
single non-aliased pointer, the pass will either remove all the uses of
the pointer or not change anything.

To help review this patch, I recommend starting by looking at the
interfaces defined in `Mem2Reg.td`, along with their reference
implementation for LLVM IR defined in `LLVMMem2Reg.cpp`. Then, the core
algorithm is implemented in `Mem2Reg.cpp`.

If this is all good I also have an implementation of the interfaces for
0-dimensional memref promotion that I can upstream afterwards.

Reviewed By: gysit

Differential Revision: https://reviews.llvm.org/D148109
2023-04-27 06:00:48 +00:00
Fabian Mora
54e96f4f97 [mlir][GPUDialect] Implement memory attributions for LaunchOp
Currently memory attributions are not supported for gpu::LaunchOp, this patch implements memory attributions for gpu::LaunchOp and modifies the KernelOutlining pass to make the attributions available in GPUFuncOp.

Reviewed By: makslevental

Differential Revision: https://reviews.llvm.org/D147809
2023-04-26 17:53:18 -05:00
max
5b303f21d3 [MLIR][python bindings] Reimplement replace_all_uses_with on PyValue
Differential Revision: https://reviews.llvm.org/D149261
2023-04-26 14:04:33 -05:00
Jan Sjodin
d3f9388ffb [OpenMP][Flang][MLIR] Add lowering of TargetOp for host codegen to LLVM-IR
This patch adds lowering of TargetOps for the host. The lowering outlines the
target region function and uses the OpenMPIRBuilder support functions to emit
the function and call. Code generation for offloading will be done in later
patches.

Reviewed By: kiranchandramohan, jdoerfert, agozillon

Differential Revision: https://reviews.llvm.org/D147172
2023-04-26 12:51:31 -04:00
Quentin Colombet
e02d4142dc [MemRefToLLVM] Add a method in MemRefDescriptor to get the buffer addr
This patch pushes the computation of the start address of a memref in one
place (a method in MemRefDescriptor.)

This allows all the (indirect) users of this method to produce the start
address in the same way.

Thanks to this change, we expose more CSEs opportunities and thanks to
that, the backend is able to properly find the `llvm.assume` expression
related to the base address as demonstrated in the added test.

Differential Revision: https://reviews.llvm.org/D148947
2023-04-26 06:12:18 +02:00
Jakub Kuderski
da0730b908 [mlir][arith] Add initial integer bitwidth narrowing pass
This pass reduces the logical complexity of arith ops by choosing
narrowest supported operand bitwidth. On some targets like mobile GPUs,
narrower bitwidths also bring better runtime performance.

The first batch of rewrites handles a simple case of `arith.sitofp`
and `arith.uitofp` with zero/sign-extended inputs. In future revisions,
I plan to extend it with the following:
-  Propagating sign/zero-extensions through bit-pattern-preserving ops,
   e.g., vector transpose, broadcast, insertions/extractions.
-  Handling `linalg.index` using the `ValueBounds` interface.
-  Handling more arith ops.

Reviewed By: springerm, antiagainst

Differential Revision: https://reviews.llvm.org/D149118
2023-04-25 22:33:11 -04:00
Quinn Dawkins
650f04feda [mlir][vector] Add pattern to break down vector.bitcast
The pattern added here is intended as a last resort for targets like
SPIR-V where there are vector size restrictions and we need to be able
to break down large vector types. Vectorizing loads/stores for small
bitwidths (e.g. i8) relies on bitcasting to a larger element type and
patterns to bubble bitcast ops to where they can cancel.
This fails for cases such as
```
%1 = arith.trunci %0 : vector<2x32xi32> to vector<2x32xi8>
vector.transfer_write %1, %destination[%c0, %c0] {in_bounds = [true, true]} : vector<2x32xi8>, memref<2x32xi8>
```
where the `arith.trunci` op essentially does the job of one of the
bitcasts, leading to a bitcast that need to be further broken down
```
vector.bitcast %0 : vector<16xi8> to vector<4xi32>
```

Differential Revision: https://reviews.llvm.org/D149065
2023-04-25 20:18:02 -04:00
Alex Zinenko
135d29c3f5 [mlir] reorgnize Linalg TransformOps files. NFC
Mirror the separation between LinalgTransformOps and LinalgMatchOps in
headers. Create a separate pair of files for the extension.

Depends on D148017

Reviewed By: springerm

Differential Revision: https://reviews.llvm.org/D148075
2023-04-25 22:44:01 +00:00
max
fd527ceff1 Revert "[MLIR][python bindings] implement replace_all_uses_with on PyValue"
This reverts commit 3bab7cb089 because it breaks sanitizers.

Differential Revision: https://reviews.llvm.org/D149188
2023-04-25 15:45:17 -05:00
Mehdi Amini
9c8db444bc Remove deprecated preloadDialectInContext flag for MlirOptMain that has been deprecated for 2 years
See https://discourse.llvm.org/t/psa-preloaddialectincontext-has-been-deprecated-for-1y-and-will-be-removed/68992

Differential Revision: https://reviews.llvm.org/D149039
2023-04-24 14:37:31 -07:00
Mehdi Amini
ffd6f6b91a Remove deprecated entry point for MlirOptMain
See: https://discourse.llvm.org/t/psa-migrating-mlir-opt-like-tools-to-use-mliroptmainconfig/68991

Differential Revision: https://reviews.llvm.org/D149038
2023-04-24 14:37:31 -07:00
Mehdi Amini
cca510640b Refactor the mlir-opt command line options related to debugging in a helper
This makes it reusable across various tooling and reduces the amount of
boilerplate needed.

Differential Revision: https://reviews.llvm.org/D144818
2023-04-24 14:34:15 -07:00
Mehdi Amini
1020150e7a Add a GDB/LLDB interface for interactive debugging of MLIR Actions
This includes a small runtime acting as callback for the ExecutionEngine
and a C API that makes it possible to control from the debugger.

A python script for LLDB is included that hook a new `mlir` subcommand
and allows to set breakpoints and inspect the current action, the context
and the stack.

Differential Revision: https://reviews.llvm.org/D144817
2023-04-24 14:34:15 -07:00
Jacques Pienaar
0911558005 [mlir] Dialect type/attr bytecode read/write generator.
Tool to help generate dialect bytecode Attribute & Type reader/writing.
Show usage by flipping builtin dialect.

It helps reduce boilerplate when writing dialect bytecode attribute and
type readers/writers. It is not an attempt at a generic spec mechanism
but rather practically focussing on boilerplate reduction while also
considering that it need not be the only in memory format and make it
relatively easy to change.

There should be some cleanup in follow up as we expand to more dialects.

Differential Revision: https://reviews.llvm.org/D144820
2023-04-24 11:53:58 -07:00
Quinn Dawkins
435f7d4c2e [mlir][vector] Add unroll pattern for vector.gather
This pattern is useful for SPIR-V to unroll to a supported vector size
before later lowerings. The unrolling pattern is closer to an
elementwise op than the transfer ops because the index values from which
to extract elements are captured by the index vector and thus there is
no need to update the base offsets when unrolling gather.

Differential Revision: https://reviews.llvm.org/D149066
2023-04-24 14:02:59 -04:00
Hanhan Wang
6f87b50be6 [mlir][linalg] Add support for lowering pack with outer_dims_perm.
Reviewed By: chelini, qcolombet

Differential Revision: https://reviews.llvm.org/D148845
2023-04-24 10:39:37 -07:00
max
98fbd9d3f9 [MLIR][python bindings] implement replace_all_uses_with on PyValue
Differential Revision: https://reviews.llvm.org/D148816
2023-04-24 10:08:43 -05:00
Andrew Gozillon
2e634367be [Flang][OpenMP][Driver][MLIR] Port fopenmp-host-ir-file-path flag and add MLIR module attribute to proliferate to OpenMP IR lowering
This patch ports the fopenmp-host-ir-file-path flag from Clang to Flang-new, this flag
is added by the driver to the device pass when doing two phase compilation (device + host).

This flag is then applied to the module when compiling during the OpenMP device phase.
This file can then be utilised during lowering of the OpenMP dialect to LLVM-IR, which
allows the device and host to maintain 1:1 mapping of OpenMP metadata for variables
during lowering via the OpenMPIRBuilders loadOffloadInfoMetadata facilities
(which is used for declare target and I believe target regions as well).

Reviewer: awarzynski

Differential Revision: https://reviews.llvm.org/D148038
2023-04-24 10:06:21 -05:00
Christian Ulmann
dc2b8ae962 [LoopInfo] SFINAE mechanism for hoist into check
This commit introduces an SFINAE mechanism to make the LLVM hoist into
check member function not leak into the users of LoopInfo.

Depends on D148235

Reviewed By: ftynse

Differential Revision: https://reviews.llvm.org/D148504
2023-04-24 06:21:17 +00:00
Christian Ulmann
f5425c128a [LoopInfo] Move generic LoopInfo into own files
This commit splits the generic part of `LoopInfo` into separate files.
These new `GenericLoopInfo` files are located in `llvm/Support` to be inline
with `GenericDomTree`.

Furthermore, this change ensures that MLIR's Bazel build does not have
to link against `LLVMAnalysis` just to use these template headers.

Depends on D148219

Reviewed By: ftynse

Differential Revision: https://reviews.llvm.org/D148235
2023-04-24 06:07:05 +00:00
Mehdi Amini
67205f907a Remove -action suffix from the tag defining an Action: it is redundant here
Differential Revision: https://reviews.llvm.org/D149036
2023-04-23 23:06:27 -06:00
Hanhan Wang
8d163e5045 [mlir][Vector] Add 16x16 strategy to vector.transpose lowering.
It adds a `shuffle_16x16` strategy LowerVectorTranspose and renames `shuffle` to `shuffle_1d`. The idea is similar to 8x8 cases in x86Vector::avx2. The general algorithm is:

```
interleave 32-bit lanes using
    8x _mm512_unpacklo_epi32
    8x _mm512_unpackhi_epi32
interleave 64-bit lanes using
    8x _mm512_unpacklo_epi64
    8x _mm512_unpackhi_epi64
permute 128-bit lanes using
   16x _mm512_shuffle_i32x4
permute 256-bit lanes using again
   16x _mm512_shuffle_i32x4
```

After the first stage, they got transposed to

```
 0  16   1  17   4  20   5  21   8  24   9  25  12  28  13  29
 2  18   3  19   6  22   7  23  10  26  11  27  14  30  15  31
32  48  33  49 ...
34  50  35  51 ...
64  80  65  81 ...
...
```

After the second stage, they got transposed to

```
 0  16  32  48 ...
 1  17  33  49 ...
 2  18  34  49 ...
 3  19  35  51 ...
64  80  96 112 ...
65  81  97 114 ...
66  82  98 113 ...
67  83  99 115 ...
...
```

After the thrid stage, they got transposed to

```
  0  16  32  48   8  24  40  56  64  80  96  112 ...
  1  17  33  49 ...
  2  18  34  50 ...
  3  19  35  51 ...
  4  20  36  52 ...
  5  21  37  53 ...
  6  22  38  54 ...
  7  23  39  55 ...
128 144 160 176 ...
129 145 161 177 ...
...
```

After the last stage, they got transposed to

```
0  16  32  48  64  80  96 112 ... 240
1  17  33  49  66  81  97 113 ... 241
2  18  34  50  67  82  98 114 ... 242
...
15  31  47  63  79  96 111 127 ... 255
```

Reviewed By: dcaballe

Differential Revision: https://reviews.llvm.org/D148685
2023-04-23 11:05:41 -07:00
Mathieu Fehr
8ac8c922fb [mlir][irdl] Add IRDL registration
This patch add support for loading IRDL dialects at runtime
with `mlir-opt`.

Given the following `dialect.irdl` file:
```mlir
module {
  irdl.dialect @cmath {
    irdl.type @complex {
      %0 = irdl.is f32
      %1 = irdl.is f64
      %2 = irdl.any_of(%0, %1)
      irdl.parameters(%2)
    }

    irdl.operation @norm {
      %0 = irdl.any
      %1 = irdl.parametric @complex<%0>
      irdl.operands(%1)
      irdl.results(%0)
    }
}
```

the IRDL file can be loaded with the `mlir-opt --irdl-file=dialect.irdl`
command, and the following file can then be parsed:

```mlir
func.func @conorm(%p: !cmath.complex<f32>, %q: !cmath.complex<f32>) -> f32 {
  %norm_p = "cmath.norm"(%p) : (!cmath.complex<f32>) -> f32
  %norm_q = "cmath.norm"(%q) : (!cmath.complex<f32>) -> f32
  %pq = arith.mulf %norm_p, %norm_q : f32
  return %pq : f32
}
```

To minimize the size of this patch, the operation, attribute, and type verifier are all always returning `success()`.

Depends on D144692

Reviewed By: rriddle, Mogball, mehdi_amini

Differential Revision: https://reviews.llvm.org/D144693
2023-04-23 17:28:44 +01:00
Ramiro Leal-Cavazos
44baa65589 Revert "Revert "Fix handling of special and large vals in expand pattern for round" and "Add pattern that expands math.roundeven into math.round + arith""
This reverts commit 87cef78fa1.

The issue in the original revert is that a lit test expecting a `-nan`
as an output was failing on M2. Since the IEEE 754-2008 standard does
not require the sign to be printed when displaying a `nan`, this
commit changes the `CHECK` for `-nan` to one that checks the result
value bitcasted to an `i32` to ensure that input is being left
unchanged. This check should now be independent of platform being used
to run test.

Reviewed By: jpienaar, mehdi_amini

Differential Revision: https://reviews.llvm.org/D148941
2023-04-22 07:15:40 -07:00
Rishabh Bali
68b05a8fca [mlir] Fix i8/f32 confusion in comment
Fixes https://github.com/llvm/llvm-project/issues/62275

Differential Revision: https://reviews.llvm.org/D148952
2023-04-22 08:30:34 +02:00
Mehdi Amini
7f069f5ef4 Add a breakpoint manager that matches based on File/Line/Col Locations
This will match the locations attached to the IRunits passed in as context
with an action.

This is a recommit of d09c80515d after fixing the test on Windows.

Differential Revision: https://reviews.llvm.org/D144815
2023-04-21 23:54:20 -06:00
Aviad Cohen
2dd396c18b [mlir] tosa.reshape - Add InferTensorType interface
When this interface is used, a call to inferReturnTypeComponents()
is generated on creation and verification of the op.

Reviewed By: jpienaar

Differential Revision: https://reviews.llvm.org/D148498
2023-04-22 08:53:07 +03:00
Mehdi Amini
6fb4c9fdc1 Revert "Add a breakpoint manager that matches based on File/Line/Col Locations"
This reverts commit d09c80515d.

This is broken on Windows
2023-04-21 23:24:05 -06:00
Mehdi Amini
d09c80515d Add a breakpoint manager that matches based on File/Line/Col Locations
This will match the locations attached to the IRunits passed in as context
with an action.

Differential Revision: https://reviews.llvm.org/D144815
2023-04-21 22:28:27 -06:00
Razvan Lupusoru
7d3f341e2d [mlir][openacc] NFC Fix acc.databounds description incorrect bullet
A dash used in description is an automatic bullet in documentation
if it is the first character of the line. The dash used in the
particular location being fixed was intended to be a separator.
Thus move it to previous line to fix incorrect bullet-ization.

Reviewed By: clementval

Differential Revision: https://reviews.llvm.org/D148968
2023-04-21 15:25:31 -07:00
Hanhan Wang
ddcc50721a [mlir][linalg] Expose lowerPack and lowerUnPack utils.
Reviewed By: qcolombet

Differential Revision: https://reviews.llvm.org/D148867
2023-04-21 15:23:16 -07:00
Razvan Lupusoru
482c1dfed3 [mlir][openacc] Add descriptions for new acc data operations
The acc.bounds operation now clarifies that it is meant to be
zero-based. I added example for both C++ and Fortran to
show the normalized representation.

I also added description for each operand in the data entry
and data exit operations.

Reviewed By: clementval

Differential Revision: https://reviews.llvm.org/D148860
2023-04-21 10:21:29 -07:00
Lei Zhang
eca7698a97 [mlir][vector] NFC: Expose castAwayContractionLeadingOneDim
This commit exposes the transformation behind the pattern.
It is useful for more targeted application on a specific op
for once.

Reviewed By: kuhar

Differential Revision: https://reviews.llvm.org/D148758
2023-04-21 09:41:14 -07:00