The previous fix from af371f9f98 only applied when using a bottom-up
traversal. The change here applies the constant preprocessing logic to the
top-down case as well. This resolves the issue with the canonicalizer pass still
reordering constants, since it uses a top-down traversal by default.
Fixes#51892
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125623
The canonicalize command-line options currently have no effect, as the pass is
reading the pass options in its constructor, before they're actually
initialized. This results in the default values of the options always being used.
The change here moves the initialization of the `GreedyRewriteConfig` out of the
constructor, so that it runs after the pass options have been parsed.
Fixes#55466
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125621
Added handling rounding behavior in 32-bits for when possible. This
avoids kernel compilation generating scalarized code on platforms where
64-bit vectors are not available.
As the 48-bit lowering requires 64-bit anyway, we added a full 64-bit
solution simplifying the old path.
Reviewed By: dcaballe, mravishankar
Differential Revision: https://reviews.llvm.org/D125583
Before this fix, the bufferization implementation made the incorrect assumption that the values yielded from the "before" region must match with the values yielded from the "after" region.
Differential Revision: https://reviews.llvm.org/D125835
This diff updates the LLVMIR dialect Fastmath flags attribute to use recently
added features of `BitEnum` attributes. Specifically, this diff uses the bit
enum "group" case to represent the `fast` value as an alias for a combination
of other values (`ninf`, `nnan`, ...), instead of using a separate integer
value. (This is in line with LLVM's fastmath flags representation.) This diff
also leverages the `printBitEnumPrimaryGroups` `tblgen` field for concise
enum printing.
The `BitEnum` features were developed for an upcoming diff that adds `fastmath`
support to the arithmetic dialect. This diff simply applies some of the relevant
new features to the LLVM dialect attribute.
Reviewed By: ftynse, Mogball
Differential Revision: https://reviews.llvm.org/D124720
Previously, GEPOp relies on `findKnownStructIndices` to check if a GEP
index should be static. The truth is, `findKnownStructIndices` can only
tell you a GEP index _might_ be indexing into a struct (which should use
a static GEP index). But GEPOp::build and GEPOp::verify are falsely
taking this information as a certain answer, which creates many false
alarms like the one depicted in
`test/Target/LLVMIR/Import/dynamic-gep-index.ll`.
The solution presented here adopts a new verification scheme: When we're
recursively checking the child element types of a struct type, instead
of checking every child types, we only check the one dictated by the
(static) GEP index value. We also combine "refinement" logics --
refine/promote struct index mlir::Value into constants -- into the very
verification process since they have lots of logics in common. The
resulting code is more concise and less brittle.
We also hide GEPOp::findKnownStructIndices since most of the
aforementioned logics are already encapsulated within GEPOp::build and
GEPOp::verify, we found little reason for findKnownStructIndices (or the
new findStructIndices) to be public.
Differential Revision: https://reviews.llvm.org/D124935
This follows the same general structure of the MLIR and PDLL language
servers. This commits adds the basic functionality for setting up the server,
and initially only supports providing diagnostics. Followon commits will
build out more comprehensive behavior.
Realistically this should eventually live in llvm/, but building in MLIR is an easier
initial step given that:
* All of the necessary LSP functionality is already here
* It allows for proving out useful language features (e.g. compilation databases)
without affecting wider scale tablegen users
* MLIR has a vscode extension that can immediately take advantage of it
Differential Revision: https://reviews.llvm.org/D125440
In the overwhelmingly majority of cases only one dialect is generated at a time
anyways, and this restriction more easily catches user error when multiple
dialects might be generated. We hit this semi-recently with the PDL dialect,
and circt+other downstream users are also actively hitting this as well.
Differential Revision: https://reviews.llvm.org/D125651
An attribute without a type builder followed by a colon in an assembly format is potentially ambiguous because the parser will read ahead to parse the colon-type and pass this as the type argument to the attribute's constructor.
However, the previous verifier that checks for this ambiguity erroneously produces an error in the case of
```
let assemblyFormat = "( `(` $attr `)` )? `:`";
```
This patch fixes the bug by implementing a checker that correctly handles all edge cases, including very strange assembly formats like:
```
let assemblyFormat = "( `(` $attr ) : (`>`)? attr-dict (`>` $a^) : (`<`)? `:`";
```
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125445
This was carry over from LLVM IR where the alias definition can
be ambiguous, but MLIR type aliases have no such problems.
Having the `type` keyword is superfluous and doesn't add anything.
This commit drops it, which also nicely aligns with the syntax for
attribute aliases (which doesn't have a keyword).
Differential Revision: https://reviews.llvm.org/D125501
This patch adds a topological sort utility and pass. A topological sort reorders
the operations in a block without SSA dominance such that, as much as possible,
users of values come after their producers.
The utility function sorts topologically the operation range in a given block
with an optional user-provided callback that can be used to virtually break cycles.
The toposort pass itself recursively sorts graph regions under the target op.
Reviewed By: mehdi_amini
Differential Revision: https://reviews.llvm.org/D125063
The attribute self type parameter is currently treated like any other attribute parameter in the assembly format. The self type parameter should be handled by the operation parser and printer and play no role in the generated parsers and printers of attributes.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125724
This is the first implementation of complex (f64 and f32) support
in the sparse compiler, with complex add/mul as first operations.
Note that various features are still TBD, such as other ops, and
reading in complex values from file. Also, note that the
std::complex<float> had a bit of an ABI issue when passed as
single argument. It is still TBD if better solutions are possible.
Reviewed By: bixia
Differential Revision: https://reviews.llvm.org/D125596
We were custom counting per bit for the clz instruction. Math dialect
now has an intrinsic to do this in one instruction. Migrated to this
instruction and fixed a minor bug math-to-llvm for the intrinsic.
Reviewed By: mravishankar
Differential Revision: https://reviews.llvm.org/D125592
This changes replaces the `fully-dynamic-layout-maps` options (which was badly named) with two new options:
* `unknown-type-conversion` controls the layout maps on buffer types for which no layout map can be inferred.
* `function-boundary-type-conversion` controls the layout maps on buffer types inside of function signatures.
Differential Revision: https://reviews.llvm.org/D125615
Erase gpu.memcpy op when only uses of dest are
the memcpy op in question, its allocation and deallocation
ops.
Reviewed By: bondhugula, csigg
Differential Revision: https://reviews.llvm.org/D124257
The warning caused build errors on a couple flang testers that are
building with -Werror. The diagnostic change makes the generated
error correct.
This is a followup to https://reviews.llvm.org/D125549
Differential Revision: https://reviews.llvm.org/D125587
There are a lot of cases where we accidentally ignored the result of some
parsing hook. Mark ParseResult as LLVM_NODISCARD just like ParseResult is.
This exposed some stuff to clean up, so do.
Differential Revision: https://reviews.llvm.org/D125549
This pass is to handle computationally complex operations like
tensor.pad which are not simply lowered to the exact same operation in
the memref dialect.
Differential Revision: https://reviews.llvm.org/D125384
Implements a floating-point sign operator (using the new semi-ring ops)
that accomodates +/-Inf and +/-NaN in consistent way.
Reviewed By: bixia
Differential Revision: https://reviews.llvm.org/D125494
We shouldn't be making assumptions about the result of llvm::getTypeName,
which may have different results for anonymous namespaces depending
on the platform.
This commit refactors the current pass manager support to allow for
operation agnostic pass managers. This allows for a series of passes
to be executed on any viable pass manager root operation, instead
of one specific operation type. Op-agnostic/generic pass managers
only allow for adding op-agnostic passes.
These types of pass managers are extremely useful when constructing
pass pipelines that can apply to many different types of operations,
e.g., the default inliner simplification pipeline. With the advent of
interface/trait passes, this support can be used to define FunctionOpInterface
pass managers, or other pass managers that effectively operate on
specific interfaces/traits/etc (see #52916 for an example).
Differential Revision: https://reviews.llvm.org/D123536
This patch references code for translating memref.reinterpret_cast ops
to add translation rules for memref.reshape ops that have a static shape
argument. Since reshape ops don't have offsets, sizes, or strides, this
patch simply sets the allocated and aligned pointers of the MemRef
descriptor.
Reviewed By: ftynse, cathyzhyi
Differential Revision: https://reviews.llvm.org/D125039
Add lowering of the vector.warp_execute_on_lane_0 into scf.if plus memory
transfer for the operands and yield values.
This also add an integration test running on GPU warp. The same tests can be
later re-used with different comment lines to tests distribution
transformations.
This is mostly from @springerm contribution.
Differential Revision: https://reviews.llvm.org/D125430
Complex nested in other types is perfectly fine, just nested structs
aren't supported. Instead of checking whether there's nesting just check
whether the struct we're dealing with is a complex number.
Differential Revision: https://reviews.llvm.org/D125381
This change integrates the BufferResultsToOutParamsPass into One-Shot Module Bufferization. This improves memory management (deallocation) when buffers are returned from a function.
Note: This currently only works with statically-sized tensors. The generated code is not very efficient yet and there are opportunities for improvment (fewer copies). By default, this new functionality is deactivated.
Differential Revision: https://reviews.llvm.org/D125376
Bufferization has an optional filter to exclude certain ops from analysis+bufferization. There were a few remaining places in the codebase where the filter was not checked.
Differential Revision: https://reviews.llvm.org/D125356
When a custom operation is unknown and does not have a dialect prefix, we currently
emit an error using the name of the operation with the default dialect prefix. This
leads to a confusing error message, especially when operations get moved between dialects.
For example, `func` was recently moved out of `builtin` and to the `func` dialect. The current
error message we get is:
```
func @foo()
^ custom op 'builtin.func' is unknown
```
This could lead users to believe that there is supposed to be a `builtin.func`,
because there used to be. This commit adds a better error message that does
not assume that the operation is supposed to be in the default dialect:
```
func @foo()
^ custom op 'func' is unknown (tried 'builtin.func' as well)
```
Differential Revision: https://reviews.llvm.org/D125351
`linalg.generic` ops have canonicalizers that either remove arguments
not used in the payload, or redundant arguments. Combine these and
enhance the canonicalization to also remove results that have no use.
This is effectively dead code elimination for Linalg ops.
Differential Revision: https://reviews.llvm.org/D123632
Using "replaceUsesOfWith" is incorrect because the same initializer value may appear multiple times.
For example, if the epilogue is needed when this loop is unrolled
```
%x:2 = scf.for ... iter_args(%arg1 = %c1, %arg2 = %c1) {
...
}
```
then both epilogue's arguments will be incorrectly renamed to use the same result index (note #1 in both cases):
```
%x_unrolled:2 = scf.for ... iter_args(%arg1 = %c1, %arg2 = %c1) {
...
}
%x_epilogue:2 = scf.for ... iter_args(%arg1 = %x_unrolled#1, %arg2 = %x_unrolled#1) {
...
}
```
This is a full audit of emitError calls, I took the opportunity
to remove extranous parens and fix a couple cases where we'd
generate multiple diagnostics for the same error.
Differential Revision: https://reviews.llvm.org/D125355
Change the parsing logic to use StringRef instead of lower level
char* logic. Also, if emitting a diagnostic on the first token
in the file, we make sure to use that position instead of the
very start of the file.
Differential Revision: https://reviews.llvm.org/D125353
Move async copy operations to NVGPU as they only exist on NV target and are
designed to match ptx semantic. This allows us to also add more fine grain
caching hint attribute to the op.
Add hint to bypass L1 and hook it up to NVVM op.
Differential Revision: https://reviews.llvm.org/D125244
The current implementation of `cloneWithNewYields` has a few issues
- It clones the loop body of the original loop to create a new
loop. This is very expensive.
- It performs `erase` operations which are incompatible when this
method is called from within a pattern rewrite. All erases need to
go through `PatternRewriter`.
To address these a new utility method `replaceLoopWithNewYields` is added
which
- moves the operations from the original loop into the new loop.
- replaces all uses of the original loop with the corresponding
results of the new loop
- use a call back to allow caller to generate the new yield values.
- the original loop is modified to just yield the basic block
arguments corresponding to the iter_args of the loop. This
represents a no-op loop. The loop itself is dead (since all its uses
are replaced), but is not removed. The caller is expected to erase
the op. Consequently, this method can be called from within a
`matchAndRewrite` method of a `PatternRewriter`.
The `cloneWithNewYields` could be replaces with
`replaceLoopWithNewYields`, but that seems to trigger a failure during
walks, potentially due to the operations being moved. That is left as
a TODO.
Differential Revision: https://reviews.llvm.org/D125147
By analogy with the NVGPU dialect, introduce an AMDGPU dialect for
AMD-specific intrinsic wrappers.
The dialect initially includes wrappers around the raw buffer intrinsics.
On AMD GPUs, a memref can be converted to a "buffer descriptor" that
allows more precise control of memory access, such as by allowing for
out of bounds loads/stores to be replaced by 0/ignored without adding
additional conditional logic, which is important for performance.
The repository currently contains a limited conversion from
transfer_read/transfer_write to Mubuf intrinsics, which are an older,
deprecated intrinsic for the same functionality.
The new amdgpu.raw_buffer_* ops allow these operations to be used
explicitly and for including metadata such as whether the target
chipset is an RDNA chip or not (which impacts the interpretation of
some bits in the buffer descriptor), while still maintaining an
MLIR-like interface.
(This change also exposes the floating-point atomic add intrinsic.)
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D122765
A typical problem with missing a token is that the missing
token is at the end of a line. The problem with this is that
the error message gets reported on the start of the following
line (which is where the next / invalid token is) which can
be confusing.
Handle this by noticing this case and backing up to the end of
the previous line.
Differential Revision: https://reviews.llvm.org/D125295
Add attribute to be able to generate the intrinsic version of async copy
generating a copy with l1 bypass. This correspond to
cp.async.cg.shared.global in ptx.
Differential Revision: https://reviews.llvm.org/D125241
There are a couple of issues with the python bindings on Windows:
- `create_symlink` requires special permissions on Windows - using `copy_if_different` instead allows the build to complete and then be usable
- the path to the `python_executable` is likely to contain spaces if python is installed in Program Files. llvm's python substitution adds extra quotes in order to account for this case, but mlir's own python substitution does not
- the location of the shared libraries is different on windows
- if the type is not specified for numpy arrays, they appear to be treated as strings
I've implemented the smallest possible changes for each of these in the patch, but I would actually prefer a slightly more comprehensive fix for the python_executable and the shared libraries.
For the python substitution, I think it makes sense to leverage the existing %python instead of adding %PYTHON and instead add a new variable for the case when preloading is needed. This would also make it clearer which tests are which and should be skipped on platforms where the preloading won't work.
For the shared libraries, I think it would make sense to pass the correct path and extension (possibly even the names) to the python script since these are known by lit and don't have to be hardcoded in the test at all.
Reviewed By: stellaraccident
Differential Revision: https://reviews.llvm.org/D125122
This patch fixed the padding size calculation for Conv2d ops when the stride > 1. It contains the changes below:
- Use addBound to add constraint for AffineApplyOp in getUpperBoundForIndex. So the result value can be mapped and retrieved later.
- Fixed the bound from AffineMinOp by adding as a closed bound. Originally the bound was added as an open upper bound, which results in the incorrect bounds when we multiply the values. For example:
```
%0 = affine.min affine_map<()[s0] -> (4, -s0 + 11)>()[iv0]
%1 = affine.apply affine_map<()[s0] -> (s0 * 2)>()[%0]
If we add the affine.min as an open bound, addBound will internally transform it into the close bound "%0 <= 3". The following sliceBounds will derive the bound of %1 as "%1 <= 6" and return the open bound "%1 < 7", while the correct bound should be "%1 <= 8".
```
- In addition to addBound, I also changed sliceBounds to support returning closed upper bound, since for the size computation, we usually care about the closed bounds.
- Change the getUpperBoundForIndex to favor constant bounds when required. The sliceBounds will return a tighter but non-constant bounds, which can't be used for padding. The constantRequired option requires getUpperBoundForIndex to get the constant bounds when possible.
Reviewed By: hanchung
Differential Revision: https://reviews.llvm.org/D124821
This patch augments the `tensor-bufferize` pass by adding a conversion
rule to translate ReshapeOp from the `tensor` dialect to the `memref`
dialect, in addition to adding a unit test to validate the translation.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D125031
rkayaith