- Implement `SubsetOpInterface`, `SubsetExtractionOpInterface`,
`SubsetInsertionOpInterface` for `vector.transfer_read` and
`vector.transfer_write`.
- Move all tensor subset hoisting test cases from `Linalg` to
`loop-invariant-subset-hoisting.mlir`. (Removing 1 duplicate test case.)
The majority of subset ops operate on hyperrectangular subsets. This
commit adds a new optional interface method
(`getAccessedHyperrectangularSlice`) that can be implemented by such
subset ops. If implemented, the other `operatesOn...` interface methods
of the `SubsetOpInterface` do not have to be implemented anymore.
The comparison logic for hyperrectangular subsets (is
disjoint/equivalent) is implemented with `ValueBoundsOpInterface`. This
makes the subset hoisting more powerful: simple cases where two
different SSA values always have the same runtime value can now be
supported.
Add a loop-invariant subset hoisting pass to `mlir/Interfaces`. This
pass hoist loop-invariant tensor subsets (subset extraction and subset
insertion ops) from loop-like ops. Extraction ops are moved before the
loop. Insertion ops are moved after the loop. The loop body operates on
newly added region iter_args (one per extraction-insertion pair).
This new pass will be improved in subsequent commits (to support more
cases/ops) and will eventually replace
`Linalg/Transforms/SubsetHoisting.cpp`. In contrast to the existing
Linalg subset hoisting, the new pass is op interface-based
(`SubsetOpInterface` and `LoopLikeOpInterface`).
There is currently an op interface for subset insertion ops
(`SubsetInsertionOpInterface`), but not for subset extraction ops. This
commit adds `SubsetExtractionOpInterface` to `mlir/Interfaces`, as well
as a common dependent op interface: `SubsetOpInterface`.
- `SubsetOpInterface` is for ops that operate on tensor subsets. It
provides interface methods to check if two subset ops operate on
equivalent or disjoint subsets. Ops that implement this interface must
implement either `SubsetExtractionOpInterface` or
`SubsetInsertionOpInterface`.
- `SubsetExtractionOpInterface` is for ops that extract from a tensor at
a subset. E.g., `tensor.extract_slice`, `tensor.gather`,
`vector.transfer_read`. Current implemented only on
`tensor.extract_slice`.
- `SubsetInsertionOpInterface` is for ops that insert into a destination
tensor at a subset. E.g., `tensor.insert_slice`,
`tensor.parallel_insert_slice`, `tensor.scatter`,
`vector.transfer_write`. Currently only implemented on
`tensor.insert_slice`, `tensor.parallel_insert_slice`.
Other changes:
- Rename `SubsetInsertionOpInterface.td` to `SubsetOpInterface.td`.
- Add helper functions to `ValueBoundsOpInterface.cpp` for checking
whether two slices are disjoint.
The new interfaces will be utilized by a new "loop-invariant subset
hoisting"
transformation. (This new transform is roughly
what `Linalg/Transforms/SubsetHoisting.cpp` is doing, but in a generic
and interface-driven way.)
`SubsetInsertionOpInterface` is an interface for ops that insert into a
destination tensor at a subset. It is currently used by the
bufferization framework to support efficient
`tensor.extract_slice/insert_slice` bufferization and to drive "empty
tensor elimination".
This commit moves the interface to `mlir/Interfaces`. This is in
preparation of adding a new "loop-invariant subset hoisting"
transformation to
`mlir/Transforms/Utils/LoopInvariantCodeMotionUtils.cpp`, which will
utilize `SubsetInsertionOpInterface`. (This new transform is roughly
what `Linalg/Transforms/SubsetHoisting.cpp` is doing, but in a generic
and interface-driven way.)
These bugs were found with the new printf long double fuzzing. The long
double inf vs nan bug was introduced when we changed to
get_explicit_exponent. The bitcast msan issue hadn't come up previously,
but isn't a real bug, just a poisoning confusion.
Due to an issue when lowering from scf to spirv as there was no
conversion pass for index to spirv, we are motivated to add a conversion
pass from the Index dialect to the SPIR-V dialect. Furthermore, we add
the new conversion patterns to the scf-to-spirv conversion.
Fixes https://github.com/llvm/llvm-project/issues/63713
---------
Co-authored-by: Jeremy Kun <jkun@google.com>
This reverts commit f681225700. That
commit changed the organization of the tests of the transform dialect
interpreter but did not take into account some tests that were added in
the meantime.
A recent commit (#69190) broke the bazel builds. Turns out that Bazel
uses symlinks for providing the test files, which the path expansion of
the module loading mechanism did not handle correctly. This PR fixes
that.
It also reorganizes the tests better: It puts all `.mlir` files that are
included by some other test into a common `include` folder. This greatly
simplifies the definition of the dependencies between the different
`.mlir` files in Bazel's `BUILD` file. The commit also adds a comment to
all included files why these aren't tested themselves direclty and uses
the `%{fs-sep}` expansion for paths more consistently. Finally, it
uncomments all but one of the tests excluded in Bazel because they seem
to run now. (The remaining one includes a file that it itself a test, so
it would have to live *in* and *outside* of the `include` folder.)
This cleans up all external entry points that will have to deal with
non-permutations, making any subsequent refactoring much more local to
the lib files.
Recent changes (https://github.com/llvm/llvm-project/pull/66930)
disabled vector transfer ops hoisting with view-like intermediate ops.
The recommended way is to fold subview ops into transfer op indices
before invoking hoisting. That would mean now we see transfer op indices
involving dynamic values, instead of static constant values before with
subview ops. Therefore hoisting won't kick in anymore. This breaks
downstream users.
To fix it, this commit enables hoisting transfer ops with dynamic
indices by using `ValueBoundsConstraintSet` to prove ranges are disjoint
in `isDisjointTransferIndices`. Given that utility is used in many
places including op folders, right now we introduce a flag to it and
only set as true for "heavy" transforms in hoisting and load-store
forwarding.
This revision provides the ability to use an arbitrary named sequence op
as
the entry point to a transform dialect strategy.
It is also a step towards better transform dialect usage in pass
pipelines
that need to preload a transform library rather thanparse it on the fly.
The interpreter itself is significantly simpler than its testing
counterpart
by avoiding payload/debug root tags and multiple shared modules.
In the process, the NamedSequenceOp::apply function is adapted to allow
it
being an entry point.
NamedSequenceOp is **not** extended to take the PossibleTopLevelTrait at
this
time, because the implementation of the trait is specific to allowing
one
top-level dangling op with a region such as SequenceOp or
AlternativesOp.
In particular, the verifier of PossibleTopLevelTrait does not allow for
an
empty body, which is necessary to declare a NamedSequenceOp that gets
linked
in separately before application.
In the future, we should dispense with the PossibleTopLevelTrait
altogether
and always enter the interpreter with a NamedSequenceOp.
Lastly, relevant TD linking utilities are moved to
TransformInterpreterUtils
and reused from there.
The `release` flag is misleading and its semantics are not well defined.
Originally this was meant to allow for different `LIBC_NAMESPACE`
depending on whether the code was considered stabled and released or
unstable. It appears that we may have a canary environment that is
neither released or dev. As a consequence we move the `LIBC_NAMESPACE`
definition to its own file and each environment can override this file
with whatever makes sense.
It's no longer possible to submit bitcode apps to the Apple App Store.
The tools
used to create xar archived bitcode sections inside MachO files have
been
discontinued. Additionally, the xar APIs have been deprecated since
macOS 12,
so this change removes unnecessary code from objdump and all
dependencies on
libxar.
This fixes rdar://116600767
This revision introduces a MapRef, which will support a future
generalization beyond permutations (e.g. block sparsity). This revision
also unifies the conversion/codegen paths for the sparse_tensor.new
operation from file (eg. the readers). Note that more unification is
planned as well as general affine dim2lvl and lvl2dim (all marked with
TODOs).
This PR introduces a new Op called `warpgroup.mma.store` to the NVGPU
dialect of MLIR. The purpose of this operation is to facilitate storing
fragmanted result(s) `nvgpu.warpgroup.accumulator` produced by
`warpgroup.mma` to the given memref.
An example of fragmentated matrix is given here :
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#wgmma-64n16-d
The `warpgroup.mma.store` does followings:
1) Takes one or more `nvgpu.warpgroup.accumulator` type (fragmented
results matrix)
2) Calculates indexes per thread in warp-group and stores the data into
give memref.
Here's an example usage:
```
// A warpgroup performs GEMM, results in fragmented matrix
%result1, %result2 = nvgpu.warpgroup.mma ...
// Stores the fragmented result to memref
nvgpu.warpgroup.mma.store [%result1, %result2], %matrixD :
!nvgpu.warpgroup.accumulator< fragmented = vector<64x128xf32>>,
!nvgpu.warpgroup.accumulator< fragmented = vector<64x128xf32>>
to memref<128x128xf32,3>
```
Implementing expm1 function for double precision based on exp function
algorithm:
- Reduced x = log2(e) * (hi + mid1 + mid2) + lo, where:
* hi is an integer
* mid1 * 2^-6 is an integer
* mid2 * 2^-12 is an integer
* |lo| < 2^-13 + 2^-30
- Then exp(x) - 1 = 2^hi * 2^mid1 * 2^mid2 * exp(lo) - 1 ~ 2^hi *
(2^mid1 * 2^mid2 * (1 + lo * P(lo)) - 2^(-hi) )
- We evaluate fast pass with P(lo) is a degree-3 Taylor polynomial of
(e^lo - 1) / lo in double precision
- If the Ziv accuracy test fails, we use degree-6 Taylor polynomial of
(e^lo - 1) / lo in double double precision
- If the Ziv accuracy test still fails, we re-evaluate everything in
128-bit precision.
Inserting clones requires a lot of assumptions to hold on the input IR, e.g., all writes to a buffer need to dominate all reads. This is not guaranteed by one-shot bufferization and isn't easy to verify, thus it could quickly lead to incorrect results that are hard to debug. This commit changes the mechanism of how an ownership indicator is materialized when there is not already a unique ownership present. Additionally, we don't create copies of returned memrefs anymore when we don't have ownership. Instead, we insert assert operations to make sure we have ownership at runtime, or otherwise report to the user that correctness could not be guaranteed.
In a previous patch, the printf writer was rewritten to use a single
writer class with a buffer and a callback hook. This patch refactors
scanf's reader to match conceptually.
Summary:
The parser class for stdio currently accepts different argument
providers. In-tree this is only used for a fuzzer test, however, the
proposed implementation of the GPU handling of printf / scanf will
require custom argument handlers. This makes the current approach of
using a preprocessor macro messier. This path proposed folding this
logic into a template instantiation. The downside to this is that
because the implementation of the parser class is placed into an
implementation file we need to manually instantiate the needed templates
which will slightly bloat binary size. Alternatively we could remove the
implementation file, or key off of the `libc` external packaging macro
so it is not present in the installed version.
Don't generate enums from the main VectorOps.td file as that
transitively includes enums from Arith.
---------
Co-authored-by: Nicolas Vasilache <ntv@google.com>
This is necessary to support deallocation of IR with gpu.launch
operations because it does not implement the RegionBranchOpInterface.
Implementing the interface would require it to support regions with
unstructured control flow and produced arguments/results.
This revision adds support for empty tensor elimination to
"bufferization.materialize_in_destination" by implementing the
`SubsetInsertionOpInterface`.
Furthermore, the One-Shot Bufferize conflict detection is improved for
"bufferization.materialize_in_destination".
Before this change, the `ByteCode` test failed on CentOS 7 with
devtoolset-9, because strings happen to be only 8 byte aligned. In
general though, strings have no alignment guarantee.
Increase resource alignment in test to 32 bytes.
Adjust test to sufficiently align buffer.
Add test to check error when buffer is insufficiently aligned.
Since ownership based buffer deallocation requires a few passes to be run in a somewhat fixed sequence, it makes sense to have a pipeline for convenience (and to reduce the number of transform ops to represent default deallocation).
This new interface allows operations to implement custom handling of ownership values and insertion of dealloc operations which is useful when an op cannot implement the interfaces supported by default by the buffer deallocation pass (e.g., because they are not exactly compatible or because there are some additional semantics to it that would render the default implementations in buffer deallocation invalid, or because no interfaces exist for this
kind of behavior and it's not worth introducing one plus a default implementation in buffer deallocation). Additionally, it can also be used to provide more efficient handling for a specific op than the interface based default
implementations can.
Add a new Buffer Deallocation pass with the intend to replace the old
one. For now it is added as a separate pass alongside in order to allow
downstream users to migrate over gradually. This new pass has the goal
of inserting fewer clone operations and supporting additional use-cases.
Please refer to the Buffer Deallocation section in the updated
Bufferization.md file for more information on how this new pass works.
This commit generalizes the special
tensor.extract_slice/tensor.insert_slice bufferization rules to tensor
subset ops.
Ops that insert a tensor into a tensor at a specified subset (e.g.,
tensor.insert_slice, tensor.scatter) can implement the
`SubsetInsertionOpInterface`.
Apart from adding a new op interface (extending the API), this change is
NFC. The only ops that currently implement the new interface are
tensor.insert_slice and tensor.parallel_insert_slice, and those ops were
are supported by One-Shot Bufferize.
Since buffer deallocation requires a few passes to be run in a somewhat fixed
sequence, it makes sense to have a pipeline for convenience (and to reduce the
number of transform ops to represent default deallocation).
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D159432
This new interface allows operations to implement custom handling of ownership
values and insertion of dealloc operations which is useful when an op cannot
implement the interfaces supported by default by the buffer deallocation pass
(e.g., because they are not exactly compatible or because there are some
additional semantics to it that would render the default implementations in
buffer deallocation invalid, or because no interfaces exist for this kind of
behavior and it's not worth introducing one plus a default implementation in
buffer deallocation). Additionally, it can also be used to provide more
efficient handling for a specific op than the interface based default
implementations can.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D158756
Add a new Buffer Deallocation pass replacing the old one with the goal of
inserting fewer clone operations and supporting additional use-cases.
Please refer to the Buffer Deallocation section in the updated
Bufferization.md file for more information on how this new pass works.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D158421
The OpenACC standard specifies an `atomic` construct in section 2.12 (of
3.3 spec), used to ensure that a specific location is accessed or
updated atomically. Four different clauses are allowed: `read`, `write`,
`update`, or `capture`. If no clause appears, it is as if `update` is
used.
The OpenMP specification defines the same clauses for `omp atomic`. The
types of expression and the clauses in the OpenACC spec match the OpenMP
spec exactly. The main difference is that the OpenMP specification is a
superset - it includes clauses for `hint` and `memory order`. It also
allows conditional expression statements. But otherwise, the expression
definition matches.
Thus, for OpenACC, we refactor and reuse the OpenMP implementation as
follows:
* The atomic operations are duplicated in OpenACC dialect. This is
preferable so that each language's semantics are precisely represented
even if specs have divergence.
* However, since semantics overlap, a common interface between the
atomic operations is being added. The semantics for the interfaces are
not generic enough to be used outside of OpenACC and OpenMP, and thus
new folders were added to hold common pieces of the two dialects.
* The atomic interfaces define common accessors (such as getting `x` or
`v`) which match the OpenMP and OpenACC specs. It also adds common
verifiers intended to be called by each dialect's operation verifier.
* The OpenMP write operation was updated to use `x` and `expr` to be
consistent with its other operations (that use naming based on spec).
The frontend lowering necessary to generate the dialect can also be
reused. This will be done in a follow up change.
`type_traits` and `utility` refer to each other in their
implementations. Also `type_traits` starts to become too big to be
manageable. This PR splits each function into individual files. FTR this
is [how libcxx handles large headers as
well](https://github.com/llvm/llvm-project/tree/main/libcxx/include/__type_traits).
The reland adds two missing functions : is_destructible_v and is_reference_v
`type_traits` and `utility` refer to each other in their
implementations. Also `type_traits` starts to become too big to be
manageable. This PR splits each function into individual files. FTR this
is [how libcxx handles large headers as
well](https://github.com/llvm/llvm-project/tree/main/libcxx/include/__type_traits).
There are two motivations for this change:
1. It considerably simplifies adding support for the realloc operation to the
new buffer deallocation pass by lowering the realloc such that no
deallocation operation is inserted and the deallocation pass itself can
insert that dealloc
2. The lowering is expressed on a higher level and thus easier to understand,
and the lowerings of the memref operations it is composed of don't have to
be duplicated in the MemRefToLLVM lowering (also see discussion in
https://reviews.llvm.org/D133424)
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D159430
Incorporated two header files directly into other since
other parts were used (and it makes it hard to find the
definitions). Removed TODOs that are less likely to be done.
Reviewed By: yinying-lisa-li
Differential Revision: https://reviews.llvm.org/D159381
Incorporated two header files directly into other since
other parts were used (and it makes it hard to find the
definitions). Removed TODOs that are less likely to be done.
Reviewed By: Peiming
Differential Revision: https://reviews.llvm.org/D159330
Functions are always callable operations and thus every operation
implementing the `FunctionOpInterface` also implements the
`CallableOpInterface`. The only exception was the FuncOp in the toy
example. To make implementation of the `FunctionOpInterface` easier,
this commit lets `FunctionOpInterface` inherit from
`CallableOpInterface` and merges some of their methods. More precisely,
the `CallableOpInterface` has methods to get the argument and result
attributes and a method to get the result types of the callable region.
These methods are always implemented the same way as their analogues in
`FunctionOpInterface` and thus this commit moves all the argument and
result attribute handling methods to the callable interface as well as
the methods to get the argument and result types. The
`FuntionOpInterface` then does not have to declare them as well, but
just inherits them from the `CallableOpInterface`.
Adding the inheritance relation also required to move the
`FunctionOpInterface` from the IR directory to the Interfaces directory
since IR should not depend on Interfaces.
Reviewed By: jpienaar, springerm
Differential Revision: https://reviews.llvm.org/D157988
The MatchTable-based GlobalISel Combiner backend is the new default. There are no in-tree users left of the old backend.
- Removed implementation of old MatchDAG-based Combiner, including tests, the backend itself and all supporting code.
- Renamed MatchTable backend to `GlobalISelCombinerEmitter.cpp` + removed "-matchtable" from its CL option.
- no need to have a verbose name as it's the only backend left now.
Reviewed By: aemerson
Differential Revision: https://reviews.llvm.org/D158710
This revision adds support for unstructured control flow to the bufferization infrastructure. In particular: regions with multiple blocks, `cf.br`, `cf.cond_br`.
Two helper templates are added to `BufferizableOpInterface.h`, which can be implemented by ops that supported unstructured control flow in their regions (e.g., `func.func`) and ops that branch to another block (e.g., `cf.br`).
A block signature is always bufferized together with the op that owns the block.
Differential Revision: https://reviews.llvm.org/D158094
Moves the lowering of `bufferization.dealloc` to memref into a separate pass,
but still registers the pattern in the conversion pass. This is helpful when
some tensor values (and thus `to_memref` or `to_tensor` operations) still
remain, e.g., when the function boundaries are not converted, or when constant
tensors are converted to memref.get_global at a later point.
However, it is still recommended to perform all bufferization before
deallocation to avoid memory leaks as all memref allocations inserted after the
deallocation pass was applied, have to be handled manually.
Note: The buffer deallocation pass assumes that memref values defined by
`bufferization.to_memref` don't return ownership and don't have to be
deallocated. `bufferization.to_tensor` operations are handled similarly to
`bufferization.clone` operations with the exception that the result value is
not handled because it's a tensor (not a memref).
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D159180
clangDriver depends on clangBasic, so clangBasic should not depend on
clangDriver, even just its header. Also remove clangBasic's dependency
on LLVMOption.
The issue can be seen through the bazel commit
d26dd681f9 which is reverted now.
Add hasFlagNoClaim and use it as we don't want to suppress
-Wunused-command-line-argument for -mexecute-only just because
-fsanitize= is specified.
This means llvm-mc should now build without depending on the target
CodeGen libraries.
Fix up a few includes in RISCV, AMDGPU, and X86 MCA to avoid transitive
deps on CodeGen.
Fixes#64166
Symlinks are poorly supported in Bazel+remote-exec. Create the links from `ld.lld`,
`ld64.lld`, `lld-link`, `wasm-ld` to `lld` using LLVM's `binary_alias` instead.
Reviewed By: MaskRay
Differential Revision: https://reviews.llvm.org/D157830
Also a new pass option `ConvertToLLVMPass` to populate only patterns from the specified dialects. This is needed because the existing test cases expect that only ops from certain dialects are lowered. (E.g., "arith-to-llvm" expects that only "arith" ops are lowered but not "func" ops.)
Differential Revision: https://reviews.llvm.org/D157627
This revision adds a `transform.apply_conversion_patterns.func.func_to_llvm` transformation.
It is unclear at this point whether this should be spelled out as a standalone transformation
or whether it should resemble `transform.apply_conversion_patterns.dialect_to_llvm "fun"`.
This is dependent on how we want to handle the type converter creation.
In particular the current implementation exhibits the fact that
`transform.apply_conversion_patterns.memref.memref_to_llvm_type_converter` was not rich enough
and did not match the LowerToLLVMOptions.
Keeping those options in sync across all the passes that lower to LLVM is very error prone.
Instead, we should have a single `to_llvm_type_converter`.
Differential Revision: https://reviews.llvm.org/D157553
Three different options can be specified:
* `bufferization.copy_tensor` (default)
* `linalg.copy`
* `none` (no copy_back)
Differential Revision: https://reviews.llvm.org/D156173
