Remove boilerplate examples and add a text at the dialect level to describe
what kind of operands the operations accept (i.e., scalar, tensor or vector).
Left a shorter sentence describing the input operands for each operation as
this redundancy is convenient when browsing the documentation using the
website.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D126648
This patch supports to convert the llvm intrinsic to the corresponding op. It still leaves some intrinsics to be handled specially.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D126639
The current translation uses the old "ugly"/"raw" form which used PDLValue for the arguments
and results. This commit updates the C++ generation to use the recently added sugar that
allows for directly using the desired types for the arguments and result of PDL functions.
In addition, this commit also properly imports the C++ class for ODS operations, constraints,
and interfaces. This allows for a much more convienent C++ API than previously granted
with the raw/low-level types.
Differential Revision: https://reviews.llvm.org/D124817
We were currently only completing on the first operand because
the completion check was outside of the parse loop.
Differential Revision: https://reviews.llvm.org/D124784
This commit adds a new PDLL specific LSP command, pdll.viewOutput, that
allows for viewing the intermediate outputs of a given PDLL file. The available
intermediate forms currently mirror those in mlir-pdll, namely: AST, MLIR, CPP.
This is extremely useful for a developer of PDLL, as it simplifies various testing,
and is also quite useful for users as they can easily view what is actually being
generated for their PDLL files.
This new command is added to the vscode client, and is available in the right
client context menu of PDLL files, or via the vscode command palette.
Differential Revision: https://reviews.llvm.org/D124783
This allows for the results of operations to be inferred in certain contexts,
and matches the support in PDL for result type inference. The main two
initial circumstances are when used as a replacement of another operation,
or when the operation being created implements InferTypeOpInterface.
Differential Revision: https://reviews.llvm.org/D124782
Python bindings for extensions of the Transform dialect are defined in separate
Python source files that can be imported on-demand, i.e., that are not imported
with the "main" transform dialect. This requires a minor addition to the
ODS-based bindings generator. This approach is consistent with the current
model for downstream projects that are expected to bundle MLIR Python bindings:
such projects can include their custom extensions into the bundle similarly to
how they include their dialects.
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D126208
Vectorization is a key transformation to achieve high performance on most
architectures. In the transform dialect, vectorization is implemented as a
parameterizable transform op. It currently applies to a scope of payload IR
delimited by some isolated-from-above op, mainly because several enabling
transformations (such as affine simplification) are needed to perform
vectorization and these transformation would apply to ops other than the "main"
computational payload op. A separate "navigation" transform op that obtains the
isolated-from-above ancestor of an op is introduced in the core transform
dialect. Even though it is currently only useful for vectorization,
isolated-from-above ops are a common anchor for transformations (usually
implemented as passes) that is likely to be reused in the future.
Depends On D126374
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D126542
Add ops to the structured transform extension of the transform dialect that
perform interchange, padding and scalarization on structured ops. Along with
tiling that is already defined, this provides a minimal set of transformations
necessary to build vectorizable code for a single structured op.
Define two helper traits: one that implements TransformOpInterface by applying
a function to each payload op independently and another that provides a simple
"functional-style" producer/consumer list of memory effects for the transform
ops.
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D126374
I don't see a point here in the lit tests here since sqrt, mul and other ops
expand as well. I just added "smoke" tests to verify that the conversion works
and does not create any illegal ops.
I will create a patch that adds a simple integration test to
mlir/test/Integration/Dialect/ComplexOps/ that will compare the values.
Differential Revision: https://reviews.llvm.org/D126539
This patch adds support for applying a relation on domain/range of a relation.
Reviewed By: arjunp, ftynse
Differential Revision: https://reviews.llvm.org/D126339
Now that analysis and bufferization are better separated, post-analysis steps are no longer needed. Users can directly interleave analysis and bufferization as needed.
Differential Revision: https://reviews.llvm.org/D126571
The semicolons were introduced in D126105 in order to correct clang-format, but I forgot this file must be compiled as C++98 rather than C++11.
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D126561
This supports the operation conversion for threadprivate directive. The
support for memref type conversion is not implemented.
Reviewed By: kiranchandramohan, shraiysh
Differential Revision: https://reviews.llvm.org/D124610
This patch fixes the following compiler error:
error: declaration of ‘mlir::LLVM::cconv::CConv mlir::LLVM::detail::CConvAttrStorage::CConv’ changes meaning of ‘CConv’ [-fpermissive]
CConv as a member variable name was shadowing CConv as an enumeration,
hence the compiler error.
Reviewed By: ftynse, alexbatashev
Differential Revision: https://reviews.llvm.org/D126530
This essentially builds an index for the parsed records and record values (fields).
This covers quite a few cases, but is limited by the currently lackluster location
tracking in tablegen. A followup will work on plumbing more locations through
tablegen, which should greatly improve what we can do here.
Differential Revision: https://reviews.llvm.org/D125443
This allows for following links to include files. This support is effectively
identical to the logic in the PDLL language server, and code is shared as
much as possible.
Differential Revision: https://reviews.llvm.org/D125442
This provides a format for externally specifying the include directories
for a source file. The format of the tablegen database is exactly the
same as that for PDLL, namely it includes the absolute source file name and
the set of include directories. The database format is shared to simplify
the infra, and also because the format itself is general enough to share. Even
if we desire to expand in the future to contain the actual compilation command,
nothing there is specific enough that we would need two different formats.
As with PDLL, support for generating the database is added to our mlir_tablegen
cmake command.
Differential Revision: https://reviews.llvm.org/D125441
This patch adds support for Calling Convention attribute in LLVM
dialect, including enums, custom syntax and import from LLVM IR.
Additionally fix import of dso_local attribute.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D126161
This is a followup to D126105 to move functions in SparseTensorUtils.cpp to match their locations in SparseTensorUtils.h
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D126106
This change makes the public API of SparseTensorUtils.cpp explicit, whereas before the publicity of these functions was only implicit. Implicit publicity is sufficient for mlir-opt to generate calls to these functions, but it's not enough to enable C/C++ code to call them directly in the usual way (i.e., without going through codegen). Thus, leaving the publicity implicit prevents development of other tools (e.g., microbenchmarks).
In addition this change also marks the functions MLIR_CRUNNERUTILS_EXPORT, which is required by the JIT under certain configurations (albeit not for anything in our test suite).
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D126105
The Transform dialect uses the side effect modeling mechanism to record the
effects of the transform ops on the mapping between Transform IR values and
Payload IR ops. Introduce a checker pass that warns if a Transform IR value is
used after it has been freed (consumed). This pass is mostly intended as a
debugging aid in addition to the verification/assertion mechanisms in the
transform interpreter. It reports all potential use-after-free situations.
The implementation makes a series of simplifying assumptions to be simple and
conservative. A more advanced implementation would rely on the data flow-like
analysis associated with a side-effect resource rather than a value, which is
currently not supported by the analysis infrastructure.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D126381
If we don't specify the result index while matching operand with the
result of certain operation, it's supposed to match all the results of
the operation with the operand. For registered op, it's easy to do that
by either indexing with number or name. For unregistered op, this commit
enables the numeric result indexing for this use case.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D126330
This commit fixes `Tensor_InsertSliceOp` `sizes` inputs/attributes
description.
Before this commit, the description says the `sizes` inputs/attributes
denote the size of the return type. But according to the
`InsertSliceOpConstantArgumentFolder` in
`lib/Dialect/Tensor/IR/TensorOps.cpp`, the `sizes` inputs/attributes
actually denote the size of the source type.
I had an off-line discussion with the authors of `TensorOps.td` and
`TensorOps.cpp`. We concluded that it was a typo in the Op description.
This commit updates the Op description to match the actual usage.
Differential Revision: https://reviews.llvm.org/D126264
Using 64-bit integer/float type in interface storage classes would
require Int64/Float64 capability, per the Vulkan spec:
```
shaderInt64 specifies whether 64-bit integers (signed and unsigned) are
supported in shader code. If this feature is not enabled, 64-bit integer
types must not be used in shader code. This also specifies whether
shader modules can declare the Int64 capability. Declaring and using
64-bit integers is enabled for all storage classes that SPIR-V allows
with the Int64 capability.
```
This is different from, say, 16-bit element types, where:
```
shaderInt16 specifies whether 16-bit integers (signed and unsigned) are
supported in shader code. If this feature is not enabled, 16-bit integer
types must not be used in shader code. This also specifies whether
shader modules can declare the Int16 capability. However, this only
enables a subset of the storage classes that SPIR-V allows for the Int16
SPIR-V capability: Declaring and using 16-bit integers in the Private,
Workgroup (for non-Block variables), and Function storage classes is
enabled, while declaring them in the interface storage classes (e.g.,
UniformConstant, Uniform, StorageBuffer, Input, Output, and
PushConstant) is not enabled.
```
Reviewed By: hanchung
Differential Revision: https://reviews.llvm.org/D126256
These attributes can carry useful information, e.g., pipelines
might use them to organize and chain patterns.
Reviewed By: hanchung
Differential Revision: https://reviews.llvm.org/D126320
This patch adds support for obtaining a set corresponding to the domain/range
of the relation.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D126326
This patch makes sure that the address dereferences to value in
omp.atomic.write operation.
Reviewed By: kiranchandramohan, peixin
Differential Revision: https://reviews.llvm.org/D126272
This suppresse annoying warning when building mlir.
```
warning: base class ‘class mlir::PassWrapper<{anonymous}::TestStatisticPass,
mlir::OperationPass<void> >’ should be explicitly initialized in the copy constructor [-Wextra]
```
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D126209
This commit breaks down diagnostic string literals so that the attribute
name and enumurator names can be shared with the stringify utility
function and the "expected ", " to be one of ", and ", " can be shared
between different enum-related diagnostic.
Differential Revision: https://reviews.llvm.org/D125938
Add lowering for cases where the reduction dimension is fully unrolled.
It is common to unroll the reduction dimension, therefore we would want
to lower the contractions to an elementwise vector op in this case.
Differential Revision: https://reviews.llvm.org/D126120
The tests actually require the target triple to match the host, rather than just having the host in the list of available targets. This change removes `llvm_has_native_target` and instead uses the `native` feature from the lit configuration.
Reviewed By: stellaraccident
Differential Revision: https://reviews.llvm.org/D126011
TOSAs depthwise_conv2d operation includes a reshape to include the implicit x1 dimension.
Reviewed By: rsuderman
Differential Revision: https://reviews.llvm.org/D126212
AsyncCopyOp lowering converted "size in elements" to "size in bytes"
assuming the element type size is at least one byte. This removes
that restriction, allowing for types such as i4 and b1 to be handled
correctly.
Differential Revision: https://reviews.llvm.org/D125838
This changes adds missing support for the i4 data type. Tests are added
to ensure proper lowering of an nvgpu.mma.sync operation targeting the
16x8x64xi4 and 16x8x32xi4 MMA variants in the NVVM dialect.
Differential Revision: https://reviews.llvm.org/D126092
Also fixes integration of the pass into One-Shot Bufferize and adds additional test cases.
BufferResultsToOutParams can be used with "identity-layout-map" and "fully-dynamic-layout-map". "infer-layout-map" is not supported.
Differential Revision: https://reviews.llvm.org/D125636
No longer pass static dim sizes as an attribute. This was redundant and required extra checks in the verifier. This change also makes the op symmetrical to memref::AllocOp.
Differential Revision: https://reviews.llvm.org/D126178
This diff modifies `mlir-tblgen` to generate Python Operation class `__init__()`
functions that use Python keyword-only arguments.
Previously, all `__init__()` function arguments were positional. Python code to
create MLIR Operations was required to provide values for ALL builder arguments,
including optional arguments (attributes and operands). Callers that did not
provide, for example, an optional attribute would be forced to provide `None`
as an argument for EACH optional attribute. Proposed changes in this diff use
`tblgen` record information (as provided by ODS) to generate keyword arguments
for:
- optional operands
- optional attributes (which includes unit attributes)
- default-valued attributes
These `__init__()` function keyword arguments have default `None` values (i.e.
the argument form is `optionalAttr=None`), allowing callers to create Operations
more easily.
Note that since optional arguments become keyword-only arguments (since they are
placed after the bare `*` argument), this diff will require ALL optional
operands and attributes to be provided using explicit keyword syntax. This may,
in the short term, break any out-of-tree Python code that provided values via
positional arguments. However, in the long term, it seems that requiring
keywords for optional arguments will be more robust to operation changes that
add arguments.
Tests were modified to reflect the updated Operation builder calling convention.
This diff partially addresses the requests made in the github issue below.
https://github.com/llvm/llvm-project/issues/54932
Reviewed By: stellaraccident, mikeurbach
Differential Revision: https://reviews.llvm.org/D124717
This patch updates asserts in IntegerRelation::isEqual and
IntegerRelation::isCompatible to allow these functions when number of
local identifiers are different. This change is done to reflect the
algorithmic changes done before this patch.
One of the ShuffleVectorOp::build functions checks if the incoming
vector operands is scalable vector by casting its type to
mlir::VectorType first. However, in some cases the operand is not
necessarily mlir::VectorType (e.g. it might be a LLVMVectorType).
This patch fixes this issue by using the dedicated
`LLVM::isScalableVectorType` function to determine if the incoming
vector is scalable vector or not.
Differential Revision: https://reviews.llvm.org/D125818
Add support for translating from llvm::Select, llvm::FNeg, and llvm::Unreachable.
This patch also cleans up (NFC) the opcode map for simple instructions and
adds `// clang-format off/on` comments to prevent those lines from being
churned by clang-format between commits.
Differential Revision: https://reviews.llvm.org/D125817
This change adds a new op `alloc_tensor` to the bufferization dialect. During bufferization, this op is always lowered to a buffer allocation (unless it is "eliminated" by a pre-processing pass). It is useful to have such an op in tensor land, because it allows users to model tensor SSA use-def chains (which drive bufferization decisions) and because tensor SSA use-def chains can be analyzed by One-Shot Bufferize, while memref values cannot.
This change also replaces all uses of linalg.init_tensor in bufferization-related code with bufferization.alloc_tensor.
linalg.init_tensor and bufferization.alloc_tensor are similar, but the purpose of the former one is just to carry a shape. It does not indicate a memory allocation.
linalg.init_tensor is not suitable for modelling SSA use-def chains for bufferization purposes, because linalg.init_tensor is marked as not having side effects (in contrast to alloc_tensor). As such, it is legal to move linalg.init_tensor ops around/CSE them/etc. This is not desirable for alloc_tensor; it represents an explicit buffer allocation while still in tensor land and such allocations should not suddenly disappear or get moved around when running the canonicalizer/CSE/etc.
BEGIN_PUBLIC
No public commit message needed for presubmit.
END_PUBLIC
Differential Revision: https://reviews.llvm.org/D126003
This changes adds the option to lower to NvGpu dialect ops during the
VectorToGPU convsersion pass. Because this transformation reuses
existing VectorToGPU logic, a seperate VectorToNvGpu conversion pass is
not created. The option `use-nvgpu` is added to the VectorToGPU pass.
When this is true, the pass will attempt to convert slices rooted at
`vector.contract` operations into `nvgpu.mma.sync` ops, and
`vector.transfer_read` ops are converted to either `nvgpu.ldmatrix` or
one or more `vector.load` operations. The specific data loaded will
depend on the thread id within a subgroup (warp). These index
calculations depend on data type and shape of the MMA op
according to the downstream PTX specification. The code for supporting
these details is separated into `NvGpuSupport.cpp|h`.
Differential Revision: https://reviews.llvm.org/D122940
For the hypothetical "a.b.c" op printed within a region that declares "a" as
the default dialect, MLIR would currently elide the "a." prefix and only print
"b.c". However, this becomes ambiguous while parsing as "b.c" may be exist as
the "c" op in the "b" dialect. If it does not, the parsing currently fails. Do
not elide the default dialect if the op name contains further dots to avoid the
ambiguity.
See https://discourse.llvm.org/t/dropping-dialect-prefix-for-ops-with-multiple-dots-in-the-name/62562
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125975
By closing over the `rank` itself rather than `this`, we save a method call on each iteration. A minor optimization, but one that adds up.
Depends On D126016
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D126019
Since these are unused, I've removed them from the configuration, so that it can be easier to read and follow.
Reviewed By: stellaraccident
Differential Revision: https://reviews.llvm.org/D125132
The approach I took was to define a dialect 'extern' attribute that a GlobalOp can take as a value to signify external linkage. I think this approach should compose well and should also work with wherever the OpaqueElements work goes in the future (since that is just another kind of attribute). I special cased the GlobalOp parser/printer for this case because it is significantly easier on the eyes.
In the discussion, Jeff Niu had proposed an alternative syntax for GlobalOp that I ended up not taking. I did try to implement it but a) I don't think it made anything easier to read in the common case, and b) it made the parsing/printing logic a lot more complicated (I think I would need a completely custom parser/printer to do it well). Please have a look at the common cases where the global type and initial value type match: I don't think how I have it is too bad. The less common cases seem ok to me.
I chose to only implement the direct, constant load op since that is non side effecting and there was still discussion pending on that.
Differential Revision: https://reviews.llvm.org/D124318
The current behaviour of `useDefaultTypePrinterParser` and `useDefaultAttributePrinterParser` is that they are set by default, but the dialect generator only generates the declarations for the parsing and printing hooks if it sees dialect types and attributes. Same goes for the definitions generated by the AttrOrTypeDef generator.
This can lead to confusing and undesirable behaviour if the dialect generator doesn't see the definitions of the attributes and types, for example, if they are sensibly separated into different files: `Dialect.td`, `Ops.td`, `Attributes.td`, and `Types.td`.
Now, these bits are unset by default. Setting them will always result in the dialect generator emitting the declarations for the parsing hooks. And if the AttrOrTypeDef generator sees it set, it will generate the default implementations.
Reviewed By: rriddle, stellaraccident
Differential Revision: https://reviews.llvm.org/D125809
Lowering through libm gives us a baseline version, even though it's not
going to be particularly fast. This is similar to what we do for some
math dialect ops.
Differential Revision: https://reviews.llvm.org/D125550
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
This patch cleans up multiple getMaybeValue functions to take an IdKind instead
of special functions.
Reviewed By: arjunp
Differential Revision: https://reviews.llvm.org/D125617
This patch changes `FlatAffineValueConstraints` to only allow attaching
values to non-local identifiers.
The reasoning for this change is:
1. Information attached to local identifiers can be lost since local identifiers
can be removed for output size optimizations.
2. There are no current use cases for attaching values to Local identifiers.
3. Attaching a value to a local identifier does not make sense since a local
identifier represents existential quantification.
This patch also adds some additional asserts to the affected functions.
Reviewed By: arjunp, bondhugula
Differential Revision: https://reviews.llvm.org/D125613
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
The support for this has been added by 946311b893
but then ignored by bc22b5c9a2.
This enables one to write generic code that can be instantiated for both
specific operation classes and the common base class without
specialization. Examples include functions that take/return ops, such
as:
```mlir
template <typename FnTy>
void applyIf(FnTy &&lambda, ...) {
for (Operation *op : ...) {
auto specific = dyn_cast<function_traits<FnTy>::template arg_t<0>>(op);
if (specific)
lambda(specific);
}
}
```
that would otherwise need to rely on template specialization to support
lambdas that take specific operations and those that take `Operation *`.
Differential Revision: https://reviews.llvm.org/D125543
Reviewed by: rriddle
This is a followup to D125431, to keep from confusing the machinery that generates diffs (since combining these two changes into one would obfuscate the changes actually made in the previous differential).
Depends On D125431
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D125432
This commit enables proper highlighting when inner statements are
outside of a constraint/pattern/etc. This shouldn't really happen in
actual code, but can happen in documentation (which uses the same
syntax grammar).
In addition to reducing code repetition, this also helps ensure that the various API functions follow the naming convention of mlir::sparse_tensor::primaryTypeFunctionSuffix (e.g., due to typos in the repetitious code).
Depends On D125428
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D125431
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
This enables the compiler to perform devirtualization. And benchmarks
indicate devirtualization can sometimes give considerable speedup.
Depends On D122061
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D125428
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
Op registration mechanism does not allow for ops with the same name to be
re-registered. This is okay to avoid name conflicts and debug
double-registration, but may be problematic for dialect extensions that may get
registered several times (unlike dialects that are deduplicated in the
registry). When registering ops through the Transform dialect extension
mechanism, check first if the ops are already registered and only complain in
the case of repeated registration with the same name but different TypeID.
Differential Revision: https://reviews.llvm.org/D125554
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
In this instance, the trailing return type does not improve readability
as it repeats what is returned in the same line.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125697
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
Instead of recomputing memref types from tensor types, try to infer them when possible. This results in more precise layout maps.
Differential Revision: https://reviews.llvm.org/D125614
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
In d4555698f8, the name of nano precision timer function has changed from `nano_time` to `nanoTime`, but benchmarks were not updated to reflect that. This change addresses the discrepancy.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D125217
This patch fixes a bug in areIdsUnique where it ignores the [start, end] range.
No test case is added since there are no use cases through IR from where it
can be tested, and it is hard to create a unittest since we do not currently
have Values in unittests.
Reviewed By: arjunp
Differential Revision: https://reviews.llvm.org/D124735
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
Due to an apparent bug in the Doxygen version <1.8.16 used to generate
documentation for MLIR, parts of the navigation (specifically, the lists
of inherited methods for classes) are unusable due to dynsections.js
missing from the output generated by Doxygen. Setting this flag makes
Doxygen always produce the file.
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
The current grammar is really crusty, only supports a handful of
cases, and is also out-of-date after various refactorings. This commit
refactors the textmate grammar to handle significantly more cases,
and now provides proper coloring for a majority of cases (including
dialect attributes, operations, types, etc.)
Differential Revision: https://reviews.llvm.org/D125458
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
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
We can simplify the code needed to implement dyn_cast/cast/isa support for MLIR operations with documented interfaces via the CastInfo structures. This will also provide an example of how to use CastInfo.
Depends on D123901
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D124963
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