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
Instead of requiring the client to compute the "isSplat" bit,
compute it internally. This makes the logic more consistent
and defines away a lot of "elements.size()==1" in the clients.
This addresses Issue #55185
Differential Revision: https://reviews.llvm.org/D125447
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
1. Call copy constructor of the base class
2. Assign value of the option directly
Reviewed By: dcaballe, rriddle
Differential Revision: https://reviews.llvm.org/D125101
D125214 split off a MLIRExecutionEngineUtils library that is used
by MLIRGPUTransforms. However, currently the entire ExecutionEngine
directory is skipped if the LLVM_NATIVE_ARCH target is not available.
Move the check for LLVM_NATIVE_ARCH, such that MLIRExecutionEngineUtils
always gets built, and only the JIT-related libraries are omitted
without native arch.
Differential Revision: https://reviews.llvm.org/D125357
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
Now that TableGen no longer relies on global Record state, we can allow
for the client to own the RecordKeeper and SourceMgr. Given that TableGen
internally still relies on the global llvm::SrcMgr, this method unfortunately
still isn't thread-safe.
Differential Revision: https://reviews.llvm.org/D125277
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
Before dump, Insetad of switching to generic form silently after
verification failure. Print some debug logs to help identify why an op
may be printed in a different way.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125136
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
This ensures that attributes such as the index bitwidth propagate
correctly to the AMDGPUToROCDL patterns.
Differential Revision: https://reviews.llvm.org/D125320
This patch updates calls to AnalysisState::getBuffer() so that we return
early with a failure if the call does not succeed.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D125251
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
Building libMLIR.so currently fails with:
> /usr/bin/ld: /tmp/ccNzulEA.ltrans39.ltrans.o: in function `(anonymous namespace)::SerializeToHsacoPass::optimizeLlvm(llvm::Module&, llvm::TargetMachine&)':
> /builddir/build/BUILD/llvm-project-15.0.0.src/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp:328: undefined reference to `mlir::makeOptimizingTransformer(unsigned int, unsigned int, llvm::TargetMachine*)'
This is because MLIRGPUTransforms depends on MLIRExecutionEngine in
61bb2e4ea8/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp (L328),
but MLIRExecutionEngine is marked as excluded from libMLIR.so.
However, this code doesn't require the full execution engine: It
only performs middle-end optimization, and does not need any of
the JIT/codegen infrastructure. As such, split off a separate
library MLIRExecutionEngineUtils, which only contains that part
and is not excluded from libMLIR.so.
Fixes https://github.com/llvm/llvm-project/issues/54242.
Differential Revision: https://reviews.llvm.org/D125214
Currently, building mlir with the python bindings enabled on Windows in Debug is broken because pybind11, python and cmake don't like to play together. This change normalizes how the three interact, so that the builds can now run and succeed.
The main issue is that python and cmake both make assumptions about which libraries are needed in a Windows build based on the flavor.
- cmake assumes that a debug (or a debug-like) flavor of the build will always require pythonX_d.lib and provides no option/hint to tell it to use a different library. cmake does find both the debug and release versions, but then uses the debug library.
- python (specifically pyconfig.h and by extension python.h) hardcodes the dependency on pythonX_d.lib or pythonX.lib depending on whether `_DEBUG` is defined. This is NOT transparent - it does not show up anywhere in the build logs until the link step fails with `pythonX_d.lib is missing` (or `pythonX.lib is missing`)
- pybind11 tries to "fix" this by implementing a workaround - unless Py_DEBUG is defined, `_DEBUG` is explicitly undefined right before including python headers. This also requires some windows headers to be included differently, so while clever, this is a non-trivial workaround.
mlir itself includes the pybind11 headers (which contain the workaround) AS WELL AS python.h, essentially always requiring both pythonX.lib and pythonX_d.lib for linking. cmake explicitly only adds one or the other, so the build fails.
This change does a couple of things:
- In the cmake files, explicitly add the release version of the python library on Windows builds regardless of flavor. Since Py_DEBUG is not defined, pybind11 will always require release and it will be satisfied
- To satisfy python as well, this change removes any explicit inclusions of Python.h on Windows instead relying on the fact that pybind11 headers will bring in what is needed
There are a few additional things that we could do but I rejected as unnecessary at this time:
- define Py_DEBUG based on the CMAKE_BUILD_TYPE - this will *mostly* work, we'd have to think through multiconfig generators like VS, but it's possible. There doesn't seem to be a need to link against debug python at the moment, so I chose not to overcomplicate the build and always default to release
- similar to above, but define Py_DEBUG based on the CMAKE_BUILD_TYPE *as well as* the presence of the debug python library (`Python3_LIBRARY_DEBUG`). Similar to above, this seems unnecessary right now. I think it's slightly better than above because most people don't actually have the debug version of python installed, so this would prevent breaks in that case.
- similar to the two above, but add a cmake variable to control the logic
- implement the pybind11 workaround directly in mlir (specifically in Interop.h) so that Python.h can still be included directly. This seems prone to error and a pain to maintain in lock step with pybind11
- reorganize how the pybind11 headers are included and place at least one of them in Interop.h directly, so that the header has all of its dependencies included as was the original intention. I decided against this because it really doesn't need pybind11 logic and it's always included after pybind11 is, so we don't necessarily need the python includes
Reviewed By: stellaraccident
Differential Revision: https://reviews.llvm.org/D125284
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
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
libm doesn't have overloads for the small types, so promote them to a
bigger type and use the f32 function.
Differential Revision: https://reviews.llvm.org/D125093
Adds missing logic in the lowering from NvGPU to NVVM to support fp32
(in an accumulator operand) and tf32 (in multiplicand operand) types.
Fixes logic in one of the helper functions for converting the result
of a mma.sync operation with multiple 8x256bit output tiles, which is
the case for f32 outputs.
Differential Revision: https://reviews.llvm.org/D124533
This was leftover from when the standard dialect was destroyed, and
when FuncOp moved to the func dialect. Now that these transitions
have settled a bit we can drop these.
Most updates were handled using a simple regex: replace `^( *)func` with `$1func.func`
Differential Revision: https://reviews.llvm.org/D124146
Ops that are created during the bufferization were not analyzed (when run with One-Shot Bufferize), and users should instead create memref ops directly.
Futhermore, this fixes an issue where an op was erased (and put on the `erasedOps` list), but subsequently a new tensor op was created at the same memory location. This op was then not bufferized. Disallowing the creation of new tensor ops simplifies the bufferization and fixes such issues.
Differential Revision: https://reviews.llvm.org/D125017
This follows the same implementation strategy as scf::ForOp and common functionality is extracted into helper functions.
This implementation works well in cases where each yielded value (from either body/condition region) is equivalent to the corresponding bbArg of the parent block. In that case, each OpResult of the loop may be aliasing with the corresponding OpOperand of the loop (and with no other OpOperand).
In the absence of said equivalence relationship, new buffer copies must be inserted, so that the aliasing OpOperand/OpResult contract of scf::WhileOp is honored. In essence, by yielding a newly allocated buffer, we can enforce the specified may-alias relationship. (Newly allocated buffers cannot alias with any OpOperands of the loop.)
Differential Revision: https://reviews.llvm.org/D124929
A large DenseElementsAttr of i1could trigger a bug in printer/parser roundtrip.
Ex. A DenseElementsAttr of i1 with 200 elements will print as Hex format of length 400 before the fix. However, when parsing the printed text, an error will be triggered. After fix, the printed length will be 50.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D122925
Although we now have semi-rings to deal with arbitrary ops,
it is still good to convey zero-preserving semantics of
ops to the sparse compiler.
Reviewed By: bixia
Differential Revision: https://reviews.llvm.org/D125043
wren romano