Summary:
On Windows, building `mlir_c_runner_utils` doesn't properly export
symbols, thus resulting in an implib not being created, which causes
an error when consuming LLVM from external projects.
Differential Revision: https://reviews.llvm.org/D75769
MSVC qualifies the Effect reference contextually depending on where the
template is instantiated, leading to compiler failures if there is a
different Effect class defined.
Summary:
New classes are added to ODS to enable specifying additional information on the arguments and results of an operation. These classes, `Arg` and `Res` allow for adding a description and a set of 'decorators' along with the constraint. This enables specifying the side effects of an operation directly on the arguments and results themselves.
Example:
```
def LoadOp : Std_Op<"load"> {
let arguments = (ins Arg<AnyMemRef, "the MemRef to load from",
[MemRead]>:$memref,
Variadic<Index>:$indices);
}
```
Differential Revision: https://reviews.llvm.org/D74440
This revision introduces the infrastructure for defining side-effects and attaching them to operations. This infrastructure allows for defining different types of side effects, that don't interact with each other, but use the same internal mechanisms. At the base of this is an interface that allows operations to specify the different effect instances that are exhibited by a specific operation instance. An effect instance is comprised of the following:
* Effect: The specific effect being applied.
For memory related effects this may be reading from memory, storing to memory, etc.
* Value: A specific value, either operand/result/region argument, the effect pertains to.
* Resource: This is a global entity that represents the domain within which the effect is being applied.
MLIR serves many different abstractions, which cover many different domains. Simple effects are may have very different context, for example writing to an in-memory buffer vs a database. This revision defines uses this infrastructure to define a set of initial MemoryEffects. The are effects that generally correspond to memory of some kind; Allocate, Free, Read, Write.
This set of memory effects will be used in follow revisions to generalize various parts of the compiler, and make others more powerful(e.g. DCE).
This infrastructure was originally proposed here:
https://groups.google.com/a/tensorflow.org/g/mlir/c/v2mNl4vFCUM
Differential Revision: https://reviews.llvm.org/D74439
Putting this up mainly for discussion on
how this should be done. I am interested in MLIR from
the Julia side and we currently have a strong preference
to dynamically linking against the LLVM shared library,
and would like to have a MLIR shared library.
This patch adds a new cmake function add_mlir_library()
which accumulates a list of targets to be compiled into
libMLIR.so. Note that not all libraries make sense to
be compiled into libMLIR.so. In particular, we want
to avoid libraries which primarily exist to support
certain tools (such as mlir-opt and mlir-cpu-runner).
Note that the resulting libMLIR.so depends on LLVM, but
does not contain any LLVM components. As a result, it
is necessary to link with libLLVM.so to avoid linkage
errors. So, libMLIR.so requires LLVM_BUILD_LLVM_DYLIB=on
FYI, Currently it appears that LLVM_LINK_LLVM_DYLIB is broken
because mlir-tblgen is linked against libLLVM.so and
and independent LLVM components.
Previous version of this patch broke depencies on TableGen
targets. This appears to be because it compiled all
libraries to OBJECT libraries (probably because cmake
is generating different target names). Avoiding object
libraries results in correct dependencies.
(updated by Stephen Neuendorffer)
Differential Revision: https://reviews.llvm.org/D73130
add_llvm_library and add_llvm_executable may need to create new targets with
appropriate dependencies. As a result, it is not sufficient in some
configurations (namely LLVM_BUILD_LLVM_DYLIB=on) to only call
add_dependencies(). Instead, the explicit TableGen dependencies must
be passed to add_llvm_library() or add_llvm_executable() using the DEPENDS
keyword.
Differential Revision: https://reviews.llvm.org/D74930
In cmake, it is redundant to have a target list under target_link_libraries()
and add_dependency(). This patch removes the redundant dependency from
add_dependency().
Differential Revision: https://reviews.llvm.org/D74929
CMake allows calling target_link_libraries() without a keyword,
but this usage is not preferred when also called with a keyword,
and has surprising behavior. This patch explicitly specifies a
keyword when using target_link_libraries().
Differential Revision: https://reviews.llvm.org/D75725
This revision adds the first intrinsic for llvm.matrix.multiply.
This uses the more general `LLVM_OneResultOp` for now since the goal is
to use the
specific Matrix builders that @fhahn has created recently.
When piped through:
```
opt -O3 -enable-matrix | llc -O3 -march=x86-64 -mcpu=skylake-avx512
```
this has been verified to generate ymm instructions.
Additional function attribute support will be needed to generate proper
zmm instructions but at least things run end to end.
Benchmarking will be provided separately with the experimental
metaprogramming
[ModelBuilder](https://github.com/google/iree/tree/master/experimental/ModelBuilder)
tool when ready.
Summary:
Paying off some technical debt in VectorOps, where I introduced a special
op for a fused accumulator into reduction to avoid some issues around
printing and parsing an optional accumulator. This CL merges the two
into one op again and does things the right way (still would be nice
to have "assemblyFormat" for optional operands though....).
Reviewers: nicolasvasilache, andydavis1, ftynse, rriddle
Reviewed By: nicolasvasilache
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D75699
Summary:
This revision removes all of the functionality related to successor operands on the core Operation class. This greatly simplifies a lot of handling of operands, as well as successors. For example, DialectConversion no longer needs a special "matchAndRewrite" for branching terminator operations.(Note, the existing method was also broken for operations with variadic successors!!)
This also enables terminator operations to define their own relationships with successor arguments, instead of the hardcoded "pass-through" behavior that exists today.
Differential Revision: https://reviews.llvm.org/D75318
This greatly simplifies the requirements for builders using this mechanism for managing variadic operands.
Differential Revision: https://reviews.llvm.org/D75317
The existing API for successor operands on operations is in the process of being removed. This revision simplifies a later one that completely removes the existing API.
Differential Revision: https://reviews.llvm.org/D75316
This attribute details the segment sizes for operand groups within the operation. This revision add support for automatically populating this attribute in the declarative parser.
Differential Revision: https://reviews.llvm.org/D75315
This interface contains the necessary components to provide the same builtin behavior that terminators have. This will be used in future revisions to remove many of the hardcoded constraints placed on successors and successor operands. The interface initially contains three methods:
```c++
// Return a set of values corresponding to the operands for successor 'index', or None if the operands do not correspond to materialized values.
Optional<OperandRange> getSuccessorOperands(unsigned index);
// Return true if this terminator can have it's successor operands erased.
bool canEraseSuccessorOperand();
// Erase the operand of a successor. This is only valid to call if 'canEraseSuccessorOperand' returns true.
void eraseSuccessorOperand(unsigned succIdx, unsigned opIdx);
```
Differential Revision: https://reviews.llvm.org/D75314
This allows for simplifying OpDefGen, as well providing specializing accessors for the different successor counts. This mirrors the existing traits for operands and results.
Differential Revision: https://reviews.llvm.org/D75313
The current setup of the GPU dialect is to model both the host and
device side codegen. For cases (like IREE) the host side modeling
might not directly fit its use case, but device-side codegen is still
valuable. First step in accessing just the device-side functionality
of the GPU dialect is to allow just creating a gpu.func operation from
a gpu.launch operation. In addition this change also "inlines"
operations into the gpu.func op at time of creation instead of this
being a later step.
Differential Revision: https://reviews.llvm.org/D75287
Summary:
This patch adds support for translation of the OpenMP barrier construct to LLVM
IR. The OpenMP IRBuilder is used for this translation. In this patch the code
for translation is added to the existing LLVM dialect translation to LLVM IR.
The patch includes code changes and a testcase.
Reviewers: jdoerfert, nicolasvasilache, ftynse, rriddle, mehdi_amini
Reviewed By: ftynse, rriddle, mehdi_amini
Differential Revision: https://reviews.llvm.org/D72962
output has zero rank.
While lowering to loops, no indices should be used in the load/store
operation if the buffer is zero-rank.
Differential Revision: https://reviews.llvm.org/D75391
Summary:
Expose expandAffineMap so that it can be used by lowerings defined outside of
MLIR core.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D75589
This commit adds timestamp query commands in Vulkan runner's
compute pipeline to gain insights into how long it takes to
run the compute shader. This commit also adds timing from CPU
side for VkQueueSubmit and vkQueueWaitIdle.
Differential Revision: https://reviews.llvm.org/D75531
Summary:
Looks like a refactor that was never completed.
This change removes some unused and ambiguous definitions.
Reviewed By: bondhugula, nicolasvasilache, rriddle
Differential Revision: https://reviews.llvm.org/D75586
This commit updates SPIR-V dialect to support integer signedness
by relaxing various checks for signless to just normal integers.
The hack for spv.Bitcast can now be removed.
Differential Revision: https://reviews.llvm.org/D75611
A previous commit added support for integer signedness in C++
IntegerType. This change introduces ODS definitions for
integer types and integer (element) attributes w.r.t. signedness.
This commit also updates various existing definitions' descriptions
to mention signless where suitable to make it more clear.
Positive and non-negative integer attributes are removed to avoid
the explosion of subclasses. Instead, one should use more atmoic
constraints together with Confined to model that. For example,
`Confined<..., [IntPositive]>`.
Differential Revision: https://reviews.llvm.org/D75610
Summary:
The order of the operations has fallen out of sync as operations have been renamed and new ones have been added.
Differential Revision: https://reviews.llvm.org/D75540
Recently introduced support for converting sequential reduction loops to
CFG of basic blocks in the Standard dialect makes it possible to perform
a staged conversion of parallel reduction loops into a similar CFG by
using sequential loops as an intermediate step. This is already the case
for parallel loops without reduction, so extend the pattern to support
an additional use case.
Differential Revision: https://reviews.llvm.org/D75599
Summary:
This adds an rsqrt op to the standard dialect, and lowers
it as 1 / sqrt to the LLVM dialect.
Differential Revision: https://reviews.llvm.org/D75353
Summary:
Previously, we would, for an empty file, print the somewhat confusing
Assertion `tok == curTok [...]' failed.
With this change, we now print
Parse error [...]: expected 'def' [...]
This only affects the parser from chapters 1-6, since the more advanced
chapter 7 parser is actually able to generate an empty module from an
empty file. Nonetheless, this commit also adds the additional check to
the chapter 7 parser, for consistency.
Differential Revision: https://reviews.llvm.org/D75534
Summary: This allows for attaching the attribute to CmpF as a proper argument, and thus enables the removal of a bunch of c++ code.
Differential Revision: https://reviews.llvm.org/D75539
For ODS generated operations enable querying whether there is a derived
attribute with a given name.
Rollforward of commit 5aa57c2 without using llvm::is_contained.
Summary:
Introduce support for converting loop.for operations with loop-carried values
to a CFG in the standard dialect. This is achieved by passing loop-carried
values as block arguments to the loop condition block. This block dominates
both the loop body and the block immediately following the loop, so the
arguments of this block are remain visible there.
Differential Revision: https://reviews.llvm.org/D75513
MLIR ExecutionEngine and derived tools (e.g., mlir-cpu-runner) would trigger an
assertion inside ORC JIT while ExecutionEngine is being destructed after a
failed linking due to a missing function definition. The reason for this is the
JIT lookup that may return an Error referring to strings stored internally by
the JIT. If the Error outlives the ExecutionEngine, it would want have a
dangling reference, which is currently caught by an assertion inside JIT thanks
to hand-rolled reference counting. Rewrap the error message into a string
before returning.
Differential Revision: https://reviews.llvm.org/D75508
Some attribute kinds are not supported as "value" attributes of
`llvm.mlir.constant` when translating to LLVM IR. We were correctly reporting
an error, but continuing the translation using an "undef" value instead,
leading to a surprising mix of error messages and output IR. Abort the
translation after the error is reported.
Differential Revision: https://reviews.llvm.org/D75450
Summary:
This revision fixes a -Wzero-length-array compile error that
caused e459596917 which reverted
78f9e5d098.
Also fixes a struct vs class mismatch that broke compilation with
-Werror for Windows that caused
57397eba7a.
This revision adds padding for 1-D Vector in the common case of x86
execution with a stadard data layout. This supports properly interfacing
codegen with arrays of e.g. `vector<9xf32>`.
Such vectors are already assumed padded to the next power of 2 by LLVM
codegen with the default x86 data layout:
```
define void @test_vector_add_1d_2_3(<3 x float>* nocapture readnone %0,
<3 x float>* nocapture readonly %1, i64 %2, i64 %3, i64 %4, <3 x float>*
nocapture readnone %5, <3 x float>* nocapture readonly %6, i64 %7, i64
%8, i64 %9, <3 x float>* nocapture readnone %10, <3 x float>* nocapture
%11, i64 %12, i64 %13, i64 %14) local_unnamed_addr {
%16 = getelementptr <3 x float>, <3 x float>* %6, i64 1
%17 = load <3 x float>, <3 x float>* %16, align 16
%18 = getelementptr <3 x float>, <3 x float>* %1, i64 1
%19 = load <3 x float>, <3 x float>* %18, align 16
%20 = fadd <3 x float> %17, %19
%21 = getelementptr <3 x float>, <3 x float>* %11, i64 1
```
The pointer addressing a `vector<3xf32>` is assumed aligned `@16`.
Similarly, the pointer addressing a `vector<65xf32>` is assumed aligned
`@512`.
This revision allows using objects such as `vector<3xf32>` properly with
the standard x86 data layout used in the JitRunner. Integration testing
is done out of tree, at the moment such testing fails without this
change.
Differential Revision: https://reviews.llvm.org/D75459
This revision adds a static `mlir_c_runner_utils_static` library
for the sole purpose of being linked into `mlir_runner_utils` on
Windows.
It was previously reported that:
```
`add_llvm_library(mlir_c_runner_utils SHARED CRunnerUtils.cpp)`
produces *only* a dll on windows, the linking of mlir_runner_utils fails
because target_link_libraries is looking for a .lib file as opposed to a
.dll file. I think this may be a case where either we need to use
LINK_LIBS or explicitly build a static lib as well, but I haven't tried
either yet.
```
This reverts commit 5aa57c2812.
The source code generated due to this ods change does not compile,
as it passes to few arguments to llvm::is_contained.
This revision adds a static `mlir_c_runner_utils_static` library
for the sole purpose of being linked into `mlir_runner_utils` on
Windows.
It was previously reported that:
```
`add_llvm_library(mlir_c_runner_utils SHARED CRunnerUtils.cpp)`
produces *only* a dll on windows, the linking of mlir_runner_utils fails
because target_link_libraries is looking for a .lib file as opposed to a
.dll file. I think this may be a case where either we need to use
LINK_LIBS or explicitly build a static lib as well, but I haven't tried
either yet.
```
Summary:
This revision fixes a -Wzero-length-array compile error that
caused e459596917 which reverted
78f9e5d098.
This revision adds padding for 1-D Vector in the common case of x86
execution with a stadard data layout. This supports properly interfacing
codegen with arrays of e.g. `vector<9xf32>`.
Such vectors are already assumed padded to the next power of 2 by LLVM
codegen with the default x86 data layout:
```
define void @test_vector_add_1d_2_3(<3 x float>* nocapture readnone %0,
<3 x float>* nocapture readonly %1, i64 %2, i64 %3, i64 %4, <3 x float>*
nocapture readnone %5, <3 x float>* nocapture readonly %6, i64 %7, i64
%8, i64 %9, <3 x float>* nocapture readnone %10, <3 x float>* nocapture
%11, i64 %12, i64 %13, i64 %14) local_unnamed_addr {
%16 = getelementptr <3 x float>, <3 x float>* %6, i64 1
%17 = load <3 x float>, <3 x float>* %16, align 16
%18 = getelementptr <3 x float>, <3 x float>* %1, i64 1
%19 = load <3 x float>, <3 x float>* %18, align 16
%20 = fadd <3 x float> %17, %19
%21 = getelementptr <3 x float>, <3 x float>* %11, i64 1
```
The pointer addressing a `vector<3xf32>` is assumed aligned `@16`.
Similarly, the pointer addressing a `vector<65xf32>` is assumed aligned
`@512`.
This revision allows using objects such as `vector<3xf32>` properly with
the standard x86 data layout used in the JitRunner. Integration testing
is done out of tree, at the moment such testing fails without this
change.
Differential Revision: https://reviews.llvm.org/D75459
as it broke the Werror build:
.../sources/llvm-project/mlir/include/mlir/ExecutionEngine/CRunnerUtils.h:85:16: error: zero size arrays are an extension [-Werror,-Wzero-length-array]
char padding[detail::nextPowerOf2<sizeof(T[Dim])>() - sizeof(T[Dim])];
^~~~~~~~~~~~~~~
This reverts commit 78f9e5d098.
Summary:
Make computeConversionSet bubble up errors from nested regions. Note
that this doesn't change top-level behavior - since the nested region
calls emitError, the error was visible before, just not surfaced as
quickly.
Differential Revision: https://reviews.llvm.org/D75369
Summary: For example, DenseElementsAttr currently does not properly round-trip unsigned integer values.
Differential Revision: https://reviews.llvm.org/D75374
Summary:
This revision adds padding for 1-D Vector in the common case of x86
execution with a stadard data layout. This supports properly interfacing
codegen with arrays of e.g. `vector<9xf32>`.
Such vectors are already assumed padded to the next power of 2 by LLVM
codegen with the default x86 data layout:
```
define void @test_vector_add_1d_2_3(<3 x float>* nocapture readnone %0,
<3 x float>* nocapture readonly %1, i64 %2, i64 %3, i64 %4, <3 x float>*
nocapture readnone %5, <3 x float>* nocapture readonly %6, i64 %7, i64
%8, i64 %9, <3 x float>* nocapture readnone %10, <3 x float>* nocapture
%11, i64 %12, i64 %13, i64 %14) local_unnamed_addr {
%16 = getelementptr <3 x float>, <3 x float>* %6, i64 1
%17 = load <3 x float>, <3 x float>* %16, align 16
%18 = getelementptr <3 x float>, <3 x float>* %1, i64 1
%19 = load <3 x float>, <3 x float>* %18, align 16
%20 = fadd <3 x float> %17, %19
%21 = getelementptr <3 x float>, <3 x float>* %11, i64 1
```
The pointer addressing a `vector<3xf32>` is assumed aligned `@16`.
Similarly, the pointer addressing a `vector<65xf32>` is assumed aligned
`@512`.
This revision allows using objects such as `vector<3xf32>` properly with
the standard x86 data layout used in the JitRunner. Integration testing
is done out of tree, at the moment such testing fails without this
change.
Reviewers: ftynse
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D75459
Summary: Added brackets to fix the loop trip count computation.
The brackets ensure the bounds are subtracted before we divide
the result by the step of the loop.
Differential Revision: https://reviews.llvm.org/D75449
Summary:
This is to ensure that the template declaration is seen before
any template specialization.
Reviewers: mravishankar, antiagainst, rriddle!
Differential Revision: https://reviews.llvm.org/D75442
A printer refactoring removed automatic newline printing in the printer
of a ModuleOp. As a consequence, mlir-opt no longer printed a newline
after the closing brace of a module, which made it hard to distinguish
when used from command line. Print the newline character explicitly in
mlir-opt.
Putting this up mainly for discussion on
how this should be done. I am interested in MLIR from
the Julia side and we currently have a strong preference
to dynamically linking against the LLVM shared library,
and would like to have a MLIR shared library.
This patch adds a new cmake function add_mlir_library()
which accumulates a list of targets to be compiled into
libMLIR.so. Note that not all libraries make sense to
be compiled into libMLIR.so. In particular, we want
to avoid libraries which primarily exist to support
certain tools (such as mlir-opt and mlir-cpu-runner).
Note that the resulting libMLIR.so depends on LLVM, but
does not contain any LLVM components. As a result, it
is necessary to link with libLLVM.so to avoid linkage
errors. So, libMLIR.so requires LLVM_BUILD_LLVM_DYLIB=on
FYI, Currently it appears that LLVM_LINK_LLVM_DYLIB is broken
because mlir-tblgen is linked against libLLVM.so and
and independent LLVM components.
Previous version of this patch broke depencies on TableGen
targets. This appears to be because it compiled all
libraries to OBJECT libraries (probably because cmake
is generating different target names). Avoiding object
libraries results in correct dependencies.
(updated by Stephen Neuendorffer)
Differential Revision: https://reviews.llvm.org/D73130
add_llvm_library and add_llvm_executable may need to create new targets with
appropriate dependencies. As a result, it is not sufficient in some
configurations (namely LLVM_BUILD_LLVM_DYLIB=on) to only call
add_dependencies(). Instead, the explicit TableGen dependencies must
be passed to add_llvm_library() or add_llvm_executable() using the DEPENDS
keyword.
Differential Revision: https://reviews.llvm.org/D74930
In cmake, it is redundant to have a target list under target_link_libraries()
and add_dependency(). This patch removes the redundant dependency from
add_dependency().
Differential Revision: https://reviews.llvm.org/D74929
When compiling libLLVM.so, add_llvm_library() manipulates the link libraries
being used. This means that when using add_llvm_library(), we need to pass
the list of libraries to be linked (using the LINK_LIBS keyword) instead of
using the standard target_link_libraries call. This is preparation for
properly dealing with creating libMLIR.so as well.
Differential Revision: https://reviews.llvm.org/D74864
This is avoid the user to shoot themselves in the foot and encounter
strange crashes that are confusing until one run with TSAN.
Differential Revision: https://reviews.llvm.org/D75399
This is in preparation for the next patch D75141. The purpose is to
provide a single place where LLVM dialect registers its ops as
legal/illegal.
Reviewers: ftynse, mravishankar, herhut
Subscribers: jholewinski, bixia, sanjoy.google, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, csigg, arpith-jacob, mgester, lucyrfox, aartbik, liufengdb, Joonsoo, llvm-commits
Differential Revision: https://reviews.llvm.org/D75140
Putting this up mainly for discussion on
how this should be done. I am interested in MLIR from
the Julia side and we currently have a strong preference
to dynamically linking against the LLVM shared library,
and would like to have a MLIR shared library.
This patch adds a new cmake function add_mlir_library()
which accumulates a list of targets to be compiled into
libMLIR.so. Note that not all libraries make sense to
be compiled into libMLIR.so. In particular, we want
to avoid libraries which primarily exist to support
certain tools (such as mlir-opt and mlir-cpu-runner).
Note that the resulting libMLIR.so depends on LLVM, but
does not contain any LLVM components. As a result, it
is necessary to link with libLLVM.so to avoid linkage
errors. So, libMLIR.so requires LLVM_BUILD_LLVM_DYLIB=on
FYI, Currently it appears that LLVM_LINK_LLVM_DYLIB is broken
because mlir-tblgen is linked against libLLVM.so and
and independent LLVM components
(updated by Stephen Neuendorffer)
Differential Revision: https://reviews.llvm.org/D73130
add_llvm_library and add_llvm_executable may need to create new targets with
appropriate dependencies. As a result, it is not sufficient in some
configurations (namely LLVM_BUILD_LLVM_DYLIB=on) to only call
add_dependencies(). Instead, the explicit TableGen dependencies must
be passed to add_llvm_library() or add_llvm_executable() using the DEPENDS
keyword.
Differential Revision: https://reviews.llvm.org/D74930
In cmake, it is redundant to have a target list under target_link_libraries()
and add_dependency(). This patch removes the redundant dependency from
add_dependency().
Differential Revision: https://reviews.llvm.org/D74929
When compiling libLLVM.so, add_llvm_library() manipulates the link libraries
being used. This means that when using add_llvm_library(), we need to pass
the list of libraries to be linked (using the LINK_LIBS keyword) instead of
using the standard target_link_libraries call. This is preparation for
properly dealing with creating libMLIR.so as well.
Differential Revision: https://reviews.llvm.org/D74864
Previously, lib/Support/JitRunner.cpp was essentially a complete application,
performing all library initialization, along with dealing with command line
arguments and actually running passes. This differs significantly from
mlir-opt and required a dependency on InitAllDialects.h. This dependency
is significant, since it requires a dependency on all of the resulting
libraries.
This patch refactors the code so that tools are responsible for library
initialization, including registering all dialects, prior to calling
JitRunnerMain. This places the concern about what dialect to support
with the end application, enabling more extensibility at the cost of
a small amount of code duplication between tools. It also fixes
BUILD_SHARED_LIBS=on.
Differential Revision: https://reviews.llvm.org/D75272
Collect a list of conversion libraries in cmake, so we don't have to
list these explicitly in most binaries.
Differential Revision: https://reviews.llvm.org/D75222
Instead of creating extra libraries we don't really need, collect a
list of all dialects and use that instead.
Differential Revision: https://reviews.llvm.org/D75221
This matches loops with a affine.min upper bound, limiting the trip
count to a constant, and rewrites them into two loops, one with constant
upper bound and one with variable upper bound. The assumption is that
the constant upper bound loop will be unrolled and vectorized, which is
preferable if this is the hot path.
Differential Revision: https://reviews.llvm.org/D75240
Summary:
* add missing comma.
* remove "having to register them here" phrasing, since register it
is what we're doing, which made the comment a bit confusing.
* remove duplicate code.
* clarify link to chapter 3, since "folder" doesn't appear in that
chapter.
Differential Revision: https://reviews.llvm.org/D75263
Summary:
* Use bold font (not monospace) for legal/illegal.
* Say a few words about operation<->dialect precedence.
* Omit duplicate code samples.
* Indent items in bullet-point list.
Differential Revision: https://reviews.llvm.org/D75262
Summary:
* Let's use "override" when we're just doing standard baseclassing.
("Specialization" makes it sound like template specialization, which
this is not.)
* CallInterfaces.td has an include guard, so #ifdef not needed anymore.
* Omit duplicate code in code samples.
* Clarify which algorithm we're talking about.
* Mention that the ShapeInference code is code a snippet that belongs to
algorithm discussed in the paragraph above it.
* Add missing definition for createShapeInferencePass.
Differential Revision: https://reviews.llvm.org/D75260
Summary:
This revision split out a new CRunnerUtils library that supports
MLIR execution on targets without a C++ runtime.
Differential Revision: https://reviews.llvm.org/D75257
Summary:
This change does not add any functionality but merely exposes existing
static functions to make the associated transformations available
outside of their testing passes.
Differential Revision: https://reviews.llvm.org/D75232
Display the list of dialects known to mlir-opt. This is useful
for ensuring that linkage has happened correctly, for instance.
Differential Revision: https://reviews.llvm.org/D74865
Summary:
AffineApplyNormalizer provides common logic for folding affine maps that appear
in affine.apply into other affine operations that use the result of said
affine.apply. In the process, affine maps of both operations are composed.
During the composition `A.compose(B)` the symbols from the map A are placed
before those of the map B in a single concatenated symbol list. However,
AffineApplyNormalizer was ordering the operands of the operation being
normalized by iteratively appending the symbols into a single list accoridng to
the operand order, regardless of whether these operands are symbols of the
current operation or of the map that is being folded into it. This could lead
to wrong order of symbols and, when the symbols were bound to constant values,
to visibly incorrect folding of constants into affine maps as reported in
PR45031. Make sure symbols operands to the current operation are always placed
before symbols coming from the folded maps.
Update the test that was exercising the incorrect folder behavior. For some
reason, the order of symbol operands was swapped in the test input compared to
the previous operations, making it easy to assume the correct maps were
produced whereas they were swapping the symbols back due to the problem
described above.
Closes https://bugs.llvm.org/show_bug.cgi?id=45031
Differential Revision: https://reviews.llvm.org/D75247
This commit handles folding spv.LogicalAnd/spv.LogicalOr when
one of the operands is constant true/false.
Differential Revision: https://reviews.llvm.org/D75195
This revision performs some basic refactoring towards more easily defining Linalg "named" ops. Such named ops form the backbone of operations that are ubiquitous in the ML application domain.
Summary: bfloat16 is stored internally as a double, so we can't direct use Type::getIntOrFloatBitWidth.
Differential Revision: https://reviews.llvm.org/D75133
Summary:
The original patch had TODOs to add support for step computations,
which this commit addresses. The computations are expressed using
affine expressions so that the affine canonicalizers can simplify
the full bound and index computations.
Also cleans up the code a little and exposes the pass in the
header file.
Differential Revision: https://reviews.llvm.org/D75052
Affine dialect already has a map+operand simplification infrastructure in
place. Plug the recently added affine.min/max operations into this
infrastructure and add a simple test. More complex behavior of the simplifier
is already tested by other ops.
Addresses https://bugs.llvm.org/show_bug.cgi?id=45008.
Differential Revision: https://reviews.llvm.org/D75058
Originally, intrinsics generator for the LLVM dialect has been producing
customized code fragments for the translation of MLIR operations to LLVM IR
intrinsics. LLVM dialect ODS now provides a generalized version of the
translation code, parameterizable with the properties of the operation.
Generate ODS that uses this version of the translation code instead of
generating a new version of it for each intrinsic.
Differential Revision: https://reviews.llvm.org/D74893
All LLVM IR intrinsics are constructed in a similar way. The ODS definition of
the LLVM dialect in MLIR also lists multiple intrinsics, many of which
reproduce the same (or similar enough) code stanza to translate the MLIR
operation into the LLVM IR intrinsic. Provide a single base class containing
parameterizable code to build LLVM IR intrinsics given their name and the lists
of overloadable operands and results. Use this class to remove (almost)
duplicate translations for intrinsics defined in LLVMOps.td.
Differential Revision: https://reviews.llvm.org/D74889
Summary:
The mapper assigns annotations to loop.parallel operations that
are compatible with the loop to gpu mapping pass. The outermost
loop uses the grid dimensions, followed by block dimensions. All
remaining loops are mapped to sequential loops.
Differential Revision: https://reviews.llvm.org/D74963
Summary:
The RFC for this op is here: https://llvm.discourse.group/t/rfc-add-std-atomic-rmw-op/489
The std.atmomic_rmw op provides a way to support read-modify-write
sequences with data race freedom. It is intended to be used in the lowering
of an upcoming affine.atomic_rmw op which can be used for reductions.
A lowering to LLVM is provided with 2 paths:
- Simple patterns: llvm.atomicrmw
- Everything else: llvm.cmpxchg
Differential Revision: https://reviews.llvm.org/D74401
Previously C++ test passes for SPIR-V were put under
test/Dialect/SPIRV. Move them to test/lib/Dialect/SPIRV
to create a better structure.
Also fixed one of the test pass to use new
PassRegistration mechanism.
Differential Revision: https://reviews.llvm.org/D75066
mlir/lib/Parser/Parser.cpp:4484:15: warning: 'parseAssignmentList' overrides a member function but is not marked 'override' [-Winconsistent-missing-override]
ParseResult parseAssignmentList(SmallVectorImpl<OperandType> &lhs,
^
mlir/include/mlir/IR/OpImplementation.h:662:3: note: overridden virtual function is here
parseAssignmentList(SmallVectorImpl<OperandType> &lhs,
^
mlir/lib/Parser/Parser.cpp:4488:12: warning: unused variable 'type' [-Wunused-variable]
Type type;
^
This exploits the fact that the iterations of parallel loops are
independent so tiling becomes just an index transformation. This pass
only tiles the innermost loop of a loop nest.
The ultimate goal is to allow vectorization of the tiled loops, but I
don't think we're there yet with the current rewriting, as the tiled
loops don't have a constant trip count.
Differential Revision: https://reviews.llvm.org/D74954
Summary:
This details the C++ format as well as the new declarative format. This has been one of the major missing pieces from the toy tutorial.
Differential Revision: https://reviews.llvm.org/D74938
This revision add support for formatting successor variables in a similar way to operands, attributes, etc.
Differential Revision: https://reviews.llvm.org/D74789
This revision add support in ODS for specifying the successors of an operation. Successors are specified via the `successors` list:
```
let successors = (successor AnySuccessor:$target, AnySuccessor:$otherTarget);
```
Differential Revision: https://reviews.llvm.org/D74783
This range is useful when an desired API expects a range or when comparing two different ranges for equality, but the underlying data is a splat. This range removes the need to explicitly construct a vector in those cases.
Differential Revision: https://reviews.llvm.org/D74683
This matches the '(print|parse)OptionalAttrDictWithKeyword' functionality provided by the assembly parser/printer.
Differential Revision: https://reviews.llvm.org/D74682
When operations have optional attributes, or optional operands(i.e. empty variadic operands), the assembly format often has an optional section to represent these arguments. This revision adds basic support for defining an "optional group" in the assembly format to support this. An optional group is defined by wrapping a set of elements in `()` followed by `?` and requires the following:
* The first element of the group must be either a literal or an operand argument.
- This is because the first element must be optionally parsable.
* There must be exactly one argument variable within the group that is marked as the anchor of the group. The anchor is the element whose presence controls whether the group should be printed/parsed. An element is marked as the anchor by adding a trailing `^`.
* The group must only contain literals, variables, and type directives.
- Any attribute variables may be used, but only optional attributes can be marked as the anchor.
- Only variadic, i.e. optional, operand arguments can be used.
- The elements of a type directive must be defined within the same optional group.
An example of this can be seen with the assembly format for ReturnOp, which has a variadic number of operands.
```
def ReturnOp : ... {
let arguments = (ins Variadic<AnyType>:$operands);
// We only print the operands+types if there are a non-zero number
// of operands.
let assemblyFormat = "attr-dict ($operands^ `:` type($operands))?";
}
```
Differential Revision: https://reviews.llvm.org/D74681
This allows for injecting type constraints that are not direct 1-1 mappings, for example when one type is equal to the element type of another. This allows for moving over several more parsers to the declarative form.
Differential Revision: https://reviews.llvm.org/D74648
Summary:
NFC - Moved StandardOps/Ops.h to a StandardOps/IR dir to better match surrounding
directories. This is to match other dialects, and prepare for moving StandardOps
related transforms in out for Transforms and into StandardOps/Transforms.
Differential Revision: https://reviews.llvm.org/D74940
This patch implements the RFCs proposed here:
https://llvm.discourse.group/t/rfc-modify-ifop-in-loop-dialect-to-yield-values/463https://llvm.discourse.group/t/rfc-adding-operands-and-results-to-loop-for/459/19.
It introduces the following changes:
- All Loop Ops region, except for ReduceOp, terminate with a YieldOp.
- YieldOp can have variadice operands that is used to return values out of IfOp and ForOp regions.
- Change IfOp and ForOp syntax and representation to define values.
- Add unit-tests and update .td documentation.
- YieldOp is a terminator to loop.for/if/parallel
- YieldOp custom parser and printer
Lowering is not supported at the moment, and will be in a follow-up PR.
Thanks.
Reviewed By: bondhugula, nicolasvasilache, rriddle
Differential Revision: https://reviews.llvm.org/D74174
Thus far IntegerType has been signless: a value of IntegerType does
not have a sign intrinsically and it's up to the specific operation
to decide how to interpret those bits. For example, std.addi does
two's complement arithmetic, and std.divis/std.diviu treats the first
bit as a sign.
This design choice was made some time ago when we did't have lots
of dialects and dialects were more rigid. Today we have much more
extensible infrastructure and different dialect may want different
modelling over integer signedness. So while we can say we want
signless integers in the standard dialect, we cannot dictate for
others. Requiring each dialect to model the signedness semantics
with another set of custom types is duplicating the functionality
everywhere, considering the fundamental role integer types play.
This CL extends the IntegerType with a signedness semantics bit.
This gives each dialect an option to opt in signedness semantics
if that's what they want and helps code sharing. The parser is
modified to recognize `si[1-9][0-9]*` and `ui[1-9][0-9]*` as
signed and unsigned integer types, respectively, leaving the
original `i[1-9][0-9]*` to continue to mean no indication over
signedness semantics. All existing dialects are not affected (yet)
as this is a feature to opt in.
More discussions can be found at:
https://groups.google.com/a/tensorflow.org/d/msg/mlir/XmkV8HOPWpo/7O4X0Nb_AQAJ
Differential Revision: https://reviews.llvm.org/D72533
This change adds some missing arithmetic and logical operators to
`TemplatedIndexedValue` for EDSC usage.
Differential Revision: https://reviews.llvm.org/D74686
Summary: DenseElementsAttr is used to store tensor data, which in some cases can become extremely large(100s of mb). In these cases it is much more efficient to format the data as a string of hex values instead.
Differential Revision: https://reviews.llvm.org/D74922
This fixes test failures caused by a change to the name of the main
dylib, now called 'main'. It also hardens the engine against potential
future changes to the name.
Summary:
This implements the last step for lowering vector.contract progressively
to LLVM IR (except for masks). Multi-dimensional reductions that remain
after expanding all parallel dimensions are lowered into into simpler
vector.contract operations until a trivial 1-dim reduction remains.
Reviewers: nicolasvasilache, andydavis1
Reviewed By: andydavis1
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74880
Summary:
The current structure suffers from several problems, but the main one is that a construction failure is impossible to debug when using the 'get' methods. This is because we only optionally emit errors, so there is no context given to the user about the problem. This revision restructures this so that errors are always emitted, and the 'get' methods simply pass in an UnknownLoc to emit to. This allows for removing usages of the more constrained "emitOptionalLoc", as well as removing the need for the context parameter.
Fixes [PR#44964](https://bugs.llvm.org/show_bug.cgi?id=44964)
Differential Revision: https://reviews.llvm.org/D74876
Summary:
Lowers all free/batch dimensions in a vector.contract progressively
into simpler vector.contract operations until a direct vector.reduction
operation is reached. Then lowers 1-D reductions into vector.reduce.
Still TBD:
multi-dimensional contractions that remain after removing all the parallel dims
Reviewers: nicolasvasilache, andydavis1, rriddle
Reviewed By: andydavis1
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74797
It replaces DenseMap output with a SmallVector and it
removes empty loop levels from the output.
Reviewed By: andydavis1, mehdi_amini
Differential Revision: https://reviews.llvm.org/D74658
Summary: DenseElementsAttr stores float values as raw bits internally, so creating attributes just to have them unwrapped is extremely inefficient.
Differential Revision: https://reviews.llvm.org/D74818
Summary:
This trait takes three arguments: lhs, rhs, transformer. It verifies that the type of 'rhs' matches the type of 'lhs' when the given 'transformer' is applied to 'lhs'. This allows for adding constraints like: "the type of 'a' must match the element type of 'b'". A followup revision will add support in the declarative parser for using these equality constraints to port more c++ parsers to the declarative form.
Differential Revision: https://reviews.llvm.org/D74647
Summary:
This could trigger an assertion due to the block argument being used by
this block's own successor operands.
Reviewers: rriddle!
Subscribers: mehdi_amini, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74583
In some dialects, attributes may have default values that may be
determined only after shape inference. For example, attributes that
are dependent on the rank of the input cannot be assigned a default
value until the rank of the tensor is inferred.
While we can set attributes without explicit setters, referring to
the attributes via accessors instead of having to use the string
interface is better for compile time verification.
The proposed patch add one method per operation attribute that let us
set its value. The code is a very small modification of the existing
getter methods.
Differential Revision: https://reviews.llvm.org/D74143
Add an initial version of mlir-vulkan-runner execution driver.
A command line utility that executes a MLIR file on the Vulkan by
translating MLIR GPU module to SPIR-V and host part to LLVM IR before
JIT-compiling and executing the latter.
Differential Revision: https://reviews.llvm.org/D72696
Summary:
the .row.col variant turns out to be the popular one, contrary to what I
thought as .row.row. Since .row.col is so prevailing (as I inspect
cuDNN's behavior), I'm going to remove the .row.row support here, which
makes the patch a little bit easier.
Reviewers: ftynse
Subscribers: jholewinski, bixia, sanjoy.google, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74655
Summary:
This revision refactors the TypeConverter class to not use inheritance to add type conversions. It instead moves to a registration based system, where conversion callbacks are added to the converter with `addConversion`. This method takes a conversion callback, which must be convertible to any of the following forms(where `T` is a class derived from `Type`:
* Optional<Type> (T type)
- This form represents a 1-1 type conversion. It should return nullptr
or `llvm::None` to signify failure. If `llvm::None` is returned, the
converter is allowed to try another conversion function to perform
the conversion.
* Optional<LogicalResult>(T type, SmallVectorImpl<Type> &results)
- This form represents a 1-N type conversion. It should return
`failure` or `llvm::None` to signify a failed conversion. If the new
set of types is empty, the type is removed and any usages of the
existing value are expected to be removed during conversion. If
`llvm::None` is returned, the converter is allowed to try another
conversion function to perform the conversion.
When attempting to convert a type, the TypeConverter walks each of the registered converters starting with the one registered most recently.
Differential Revision: https://reviews.llvm.org/D74584
Fixing a bug where using a zero-rank shaped type operand to
linalg.generic ops hit an unrelated assert. This also meant that
lowering the operation to loops was not supported. Adding roundtrip
tests and lowering to loops test for zero-rank shaped type operand
with fixes to make the test pass.
Differential Revision: https://reviews.llvm.org/D74638
Summary: This class wraps around the various different ways to construct a range of Type, without forcing the materialization of that range into a contiguous vector.
Differential Revision: https://reviews.llvm.org/D74646
Summary:
Refer folks to the main website and make it explicit that the rendered
output is what is of interest and that the GitHub viewing experience may
not match (even though we are trying to keep it as close as possible, the
renderers do differ).
Differential Revision: https://reviews.llvm.org/D74739
The existing name is an artifact dating back to the times when we did not have
a dedicated TypeConverter infrastructure. It is also confusing with with the
name of classes using it.
Differential revision: https://reviews.llvm.org/D74707
We have one title in every doc which corresponds to `#`, in the some
there are multiple and it is expected to be h1 headers (visual elements
rather than organizational). Indent every nesting by one in all of the
docs with multiple titles.
Also fixing trailing whitespace.
Summary:
Linalg's promotion pass was only supporting f32 buffers due to how the
zero value was build for the `fill` operation.
Moreover, `promoteSubViewOperands` was returning a vector with one entry
per float subview while omitting integer subviews. For a program
with only integer subviews the return vector would be of size 0.
However, `promoteSubViewsOperands` would try to access a non zero
number of entries of this vector, resulting in a sefgault.
Reviewers: nicolasvasilache, ftynse
Reviewed By: ftynse
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74532
Summary: getEncodedSourceLocation can be very costly to compute, especially if the input line becomes very long. This revision inlines some of the verification of a few `getChecked` methods to avoid the materialization of an encoded source location.
Differential Revision: https://reviews.llvm.org/D74587
This patch extends affine data copy optimization utility with an
optional memref filter argument. When the memref filter is used, data
copy optimization will only generate copies for such a memref.
Note: this patch is just porting the memref filter feature from Uday's
'hop' branch: https://github.com/bondhugula/llvm-project/tree/hop.
Reviewed By: bondhugula
Differential Revision: https://reviews.llvm.org/D74342
The goal of this patch is to maximize CPU utilization on multi-socket or high core count systems, so that parallel computations such as LLD/ThinLTO can use all hardware threads in the system. Before this patch, on Windows, a maximum of 64 hardware threads could be used at most, in some cases dispatched only on one CPU socket.
== Background ==
Windows doesn't have a flat cpu_set_t like Linux. Instead, it projects hardware CPUs (or NUMA nodes) to applications through a concept of "processor groups". A "processor" is the smallest unit of execution on a CPU, that is, an hyper-thread if SMT is active; a core otherwise. There's a limit of 32-bit processors on older 32-bit versions of Windows, which later was raised to 64-processors with 64-bit versions of Windows. This limit comes from the affinity mask, which historically is represented by the sizeof(void*). Consequently, the concept of "processor groups" was introduced for dealing with systems with more than 64 hyper-threads.
By default, the Windows OS assigns only one "processor group" to each starting application, in a round-robin manner. If the application wants to use more processors, it needs to programmatically enable it, by assigning threads to other "processor groups". This also means that affinity cannot cross "processor group" boundaries; one can only specify a "preferred" group on start-up, but the application is free to allocate more groups if it wants to.
This creates a peculiar situation, where newer CPUs like the AMD EPYC 7702P (64-cores, 128-hyperthreads) are projected by the OS as two (2) "processor groups". This means that by default, an application can only use half of the cores. This situation could only get worse in the years to come, as dies with more cores will appear on the market.
== The problem ==
The heavyweight_hardware_concurrency() API was introduced so that only *one hardware thread per core* was used. Once that API returns, that original intention is lost, only the number of threads is retained. Consider a situation, on Windows, where the system has 2 CPU sockets, 18 cores each, each core having 2 hyper-threads, for a total of 72 hyper-threads. Both heavyweight_hardware_concurrency() and hardware_concurrency() currently return 36, because on Windows they are simply wrappers over std:🧵:hardware_concurrency() -- which can only return processors from the current "processor group".
== The changes in this patch ==
To solve this situation, we capture (and retain) the initial intention until the point of usage, through a new ThreadPoolStrategy class. The number of threads to use is deferred as late as possible, until the moment where the std::threads are created (ThreadPool in the case of ThinLTO).
When using hardware_concurrency(), setting ThreadCount to 0 now means to use all the possible hardware CPU (SMT) threads. Providing a ThreadCount above to the maximum number of threads will have no effect, the maximum will be used instead.
The heavyweight_hardware_concurrency() is similar to hardware_concurrency(), except that only one thread per hardware *core* will be used.
When LLVM_ENABLE_THREADS is OFF, the threading APIs will always return 1, to ensure any caller loops will be exercised at least once.
Differential Revision: https://reviews.llvm.org/D71775
This pass would currently build, but fail to run when this backend isn't
linked in. On the other hand, we'd like it to initialize only the NVPTX
backend, which isn't possible if we continue to build it without the
backend available. Instead of building a broken configuration, let's
skip building the pass entirely.
Differential Revision: https://reviews.llvm.org/D74592
The commit switching the calling convention for memrefs (5a1778057)
inadvertently introduced a bug in the function argument attribute conversion:
due to incorrect indexing of function arguments it was not assigning the
attributes to the arguments beyond those generated from the first original
argument. This was not caught in the commit since the test suite does have a
test for converting multi-argument functions with argument attributes. Fix the
bug and add relevant tests.
Summary: This revision adds support to the declarative parser for formatting enum attributes in the symbolized form. It uses this new functionality to port several of the SPIRV parsers over to the declarative form.
Differential Revision: https://reviews.llvm.org/D74525
Summary:
This sets the basic framework for lowering vector.contract progressively
into simpler vector.contract operations until a direct vector.reduction
operation is reached. More details will be filled out progressively as well.
Reviewers: nicolasvasilache
Reviewed By: nicolasvasilache
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74520
Implement a pass to convert gpu.launch_func op into a sequence of
Vulkan runtime calls. The Vulkan runtime API surface is huge so currently we
don't expose separate external functions in IR for each of them, instead we
expose a few external functions to wrapper libraries which manages
Vulkan runtime.
Differential Revision: https://reviews.llvm.org/D74549
Summary:
To unblock other work, this implements basic lowering based on mapping
attributes that have to be provided on all loop.parallel. The lowering
does not yet support reduce.
Differential Revision: https://reviews.llvm.org/D73893
Summary:
On some platforms the build fails "std::function is not found". The include is used in
PassManager::IRPrinterConfig::enableIRPrinting.
Differential Revision: https://reviews.llvm.org/D74469
In code generators, one can automate the translation of typed ArrayAttrs
if element attribute translators are already implemented. However, the
type of the element attribute is lost at the construction of
TypedArrayAttrBase. With this change one can inspect the element type
and generate the translation logic automatically, which will reduce the
code repetition.
Differential Revision: https://reviews.llvm.org/D73579
Summary:
As discussed in https://llvm.discourse.group/t/rfc-add-affine-parallel/350, this is the first in a series of patches to bring in support for the `affine.parallel` operation.
This first patch adds the IR representation along with custom printer/parser implementations.
Reviewers: bondhugula, herhut, mehdi_amini, nicolasvasilache, rriddle, earhart, jbruestle
Reviewed By: bondhugula, nicolasvasilache, rriddle, earhart, jbruestle
Subscribers: jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, aartbik, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74288
This patch adapts the method MemRefDescriptor::fromStaticShape to
support static non-zero offsets. The updated method uses the
getStridesAndOffset method to extract strides and offset. The patch also
adapts the test cases since sizes and strides are now set in forward
instead of reverse order.
Differential Revision: https://reviews.llvm.org/D74474
This revision prepares the ground for declaratively defining Linalg "named" ops.
Such named ops form the backbone of operations that are ubiquitous in the ML
application domain.
This revision closely related to the definition of a "Tensor Computation
Primitives Dialect" and demonstrates that ops can be expressed as declarative
configurations of the `linalg.generic` op.
Differential Revision: https://reviews.llvm.org/D74491
Summary:
This revision allows model builder to create a linalg_matmul whose body
is a vector.contract. This shows the abstractions compose nicely.
Differential Revision: https://reviews.llvm.org/D74457
Summary: This was a missed case when ValueRange was originally added, and allows for constructing a ValueRange from the arguments of a block.
Differential Revision: https://reviews.llvm.org/D74363
A memref_cast casting to a memref with a non identity map can't be
lowered to llvm. Take the following case:
```
func @invalid_memref_cast(%arg0: memref<?x?xf64>) {
%c1 = constant 1 : index
%c0 = constant 0 : index
%5 = memref_cast %arg0 : memref<?x?xf64> to memref<?x?xf64, #map1>
%25 = std.subview %5[%c0, %c0][%c1, %c1][] : memref<?x?xf64, #map1> to memref<?x?xf64, #map1>
return
}
```
When lowering the subview mlir was assuming `%5` to have an llvm type
(which is not the case as mlir failed to lower the memref_cast).
Differential Revision: https://reviews.llvm.org/D74466
Summary:
This was broken recently when moving from dialect registration via
static initializers to explicit intialization.
Differential Revision: https://reviews.llvm.org/D74480
Thus far we have been using builtin func op to model SPIR-V functions.
It was because builtin func op used to have special treatment in
various parts of the core codebase (e.g., pass pipelines, etc.) and
it's easy to bootstrap the development of the SPIR-V dialect. But
nowadays with general op concepts and region support we don't have
such limitations and it's time to tighten the SPIR-V dialect for
completeness.
This commits introduces a spv.func op to properly model SPIR-V
functions. Compared to builtin func op, it can provide the following
benefits:
* We can control the full op so we can integrate SPIR-V information
bits (e.g., function control) in a more integrated way and define
our own assembly form and enforcing better verification.
* We can have a better dialect and library boundary. At the current
moment only functions are modelled with an external op. With this
change, all ops modelling SPIR-V concpets will be spv.* ops and
registered to the SPIR-V dialect.
* We don't need to special-case func op anymore when creating
ConversionTarget declaring SPIR-V dialect as legal. This is quite
important given we'll see more and more conversions in the future.
In the process, bumps a few FuncOp methods to the FunctionLike trait.
Differential Revision: https://reviews.llvm.org/D74226
In the previous state, we were relying on forcing the linker to include
all libraries in the final binary and the global initializer to self-register
every piece of the system. This change help moving away from this model, and
allow users to compose pieces more freely. The current change is only "fixing"
the dialect registration and avoiding relying on "whole link" for the passes.
The translation is still relying on the global registry, and some refactoring
is needed to make this all more convenient.
Differential Revision: https://reviews.llvm.org/D74461
* Rename CMake target MLIROptMain to MLIROptLib:
The target provides the main library
* Rename CMake target MLIRMlirOptLib to MLIRMlirOptMain:
The target provides the main() entry function
At the moment, the Bazel configuration of TenorFlow maps the target
MlirOptLib to "lib/Support/MlirOptMain.cpp" and MlirOptMain to
"tools/mlir-opt/mlir-opt.cpp". This is the other way around in the CMake
configuration. As discussed in the context of the pull request
https://github.com/tensorflow/tensorflow/pull/36301, it seems useful to
revise the naming in the MLIR repo.
Differential Revision: https://reviews.llvm.org/D73778
Summary:
Adds affine loop fusion transformation function to LoopFusionUtils.
Updates TestLoopFusion utility to run loop fusion transformation until a fixed point is reached.
Adds unit tests to test the transformation.
Includes ASAN bug fix for D73190.
Reviewers: bondhugula, dcaballe
Reviewed By: bondhugula, dcaballe
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, aartbik, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74330
Follow-up on D72802. Turn -convert-std-to-llvm-use-alloca and
-convert-std-to-llvm-bare-ptr-memref-call-conv into pass flags
of LLVMLoweringPass.
Reviewed By: mehdi_amini
Differential Revision: https://reviews.llvm.org/D73912
Defines a tablegen class RankedIntElementsAttr. This is an integer
version of RankedFloatElementsAttr.
Differential Revision: https://reviews.llvm.org/D73764
Summary:
The lowering to NVVM and ROCm handles tanh operations differently by
mapping them to NVVM/ROCm specific intrinsics. This conflicts with
the lowering to LLVM, which uses the default llvm intrinsic. This change
declares the LLVM intrinsics to be illegal, hence disallowing the
correspondign rewrite.
Differential Revision: https://reviews.llvm.org/D74389
We have spv.entry_point_abi for specifying the local workgroup size.
It should be decorated onto input gpu.func ops to drive the SPIR-V
CodeGen to generate the proper SPIR-V module execution mode. Compared
to using command-line options for specifying the configuration, using
attributes also has the benefits that 1) we are now able to use
different local workgroup for different entry points and 2) the
tests contains the configuration directly.
Differential Revision: https://reviews.llvm.org/D74012
Summary:
After D72555 has been landed, `linalg.indexed_generic` also accepts ranked
tensor as input and output. Add a test for it.
Differential Revision: https://reviews.llvm.org/D74267
Summary:
This revision adds EDSC support for VectorOps to enable the creation of a `vector_matmul` declaratively. The `vector_matmul` is a simple configuration
of the `vector.contract` op that follows the StructuredOps abstraction.
Differential Revision: https://reviews.llvm.org/D74284
This CL refactors EDSCs to layer them better and break unnecessary
dependencies. After this refactoring, the top-level EDSC target only
depends on IR but not on Dialects anymore and each dialect has its
own EDSC directory.
This simplifies the layering and breaks cyclic dependencies.
In particular, the declarative builder + folder are made explicit and
are now confined to Linalg.
As the refactoring occurred, certain classes and abstractions that were not
paying for themselves have been removed.
Differential Revision: https://reviews.llvm.org/D74302
The current standard to llvm conversion pass lowers subview ops only if
dynamic offsets are provided. This commit extends the lowering with a
code path that uses the constant offset of the target memref for the
subview op lowering (see Example 3 of the subview op definition for an
example) if no dynamic offsets are provided.
Differential Revision: https://reviews.llvm.org/D74280
The existing (default) calling convention for memrefs in standard-to-LLVM
conversion was motivated by interfacing with LLVM IR produced from C sources.
In particular, it passes a pointer to the memref descriptor structure when
calling the function. Therefore, the descriptor is allocated on stack before
the call. This convention leads to several problems. PR44644 indicates a
problem with stack exhaustion when calling functions with memref-typed
arguments in a loop. Allocating outside of the loop may lead to concurrent
access problems in case the loop is parallel. When targeting GPUs, the contents
of the stack-allocated memory for the descriptor (passed by pointer) needs to
be explicitly copied to the device. Using an aggregate type makes it impossible
to attach pointer-specific argument attributes pertaining to alignment and
aliasing in the LLVM dialect.
Change the default calling convention for memrefs in standard-to-LLVM
conversion to transform a memref into a list of arguments, each of primitive
type, that are comprised in the memref descriptor. This avoids stack allocation
for ranked memrefs (and thus stack exhaustion and potential concurrent access
problems) and simplifies the device function invocation on GPUs.
Provide an option in the standard-to-LLVM conversion to generate auxiliary
wrapper function with the same interface as the previous calling convention,
compatible with LLVM IR porduced from C sources. These auxiliary functions
pack the individual values into a descriptor structure or unpack it. They also
handle descriptor stack allocation if necessary, serving as an allocation
scope: the memory reserved by `alloca` will be freed on exiting the auxiliary
function.
The effect of this change on MLIR-generated only LLVM IR is minimal. When
interfacing MLIR-generated LLVM IR with C-generated LLVM IR, the integration
only needs to require auxiliary functions and change the function name to call
the wrapper function instead of the original function.
This also opens the door to forwarding aliasing and alignment information from
memrefs to LLVM IR pointers in the standrd-to-LLVM conversion.
The existing lowering of gpu.block_dim added a global variable with
the WorkGroupSize decoration. This raises an error within
Vulkan/SPIR-V validation since Vulkan requires this to have a constant
initializer. This is not yet supported in SPIR-V dialect. Changing the
lowering to return the workgroup size as a constant value instead,
obtained from spv.entry_point_abi attribute gets around the issue for
now. The validation goes through since the workgroup size is specified
using spv.execution_mode operation.
This revision adds support in the declarative assembly form for printing attributes with buildable types without the type, and moves several more parsers over to the declarative form.
Differential Revision: https://reviews.llvm.org/D74276
Summary:
This revision adds a utility to generate debug locations from the IR during compilation, by snapshotting to a output stream and using the locations that operations were dumped in that stream. The new locations may either;
* Replace the original location of the operation.
old:
loc("original_source.cpp":1:1)
new:
loc("snapshot_source.mlir":10:10)
* Fuse with the original locations as NamedLocs with a specific tag.
old:
loc("original_source.cpp":1:1)
new:
loc(fused["original_source.cpp":1:1, "snapshot"("snapshot_source.mlir":10:10)])
This feature may be used by a debugger to display the code at various different levels of the IR. It would also be able to show the different levels of IR attached to a specific source line in the original source file.
This feature may also be used to generate locations for operations generated during compilation, that don't necessarily have a user source location to attach to.
This requires changes in the printer to track the locations of operations emitted in the stream. Moving forward we need to properly(and efficiently) track the number of newlines emitted to the stream during printing.
Differential Revision: https://reviews.llvm.org/D74019
Summary: This is the most common operation performed on a CallOpInterface. This just moves the existing functionality from the CallGraph so that other users can access it.
Differential Revision: https://reviews.llvm.org/D74250
Summary: This document provides insight on the rationale and the design of Symbols in MLIR, and why they are necessary.
Differential Revision: https://reviews.llvm.org/D73590
Summary:
This revision adds support for printing pass options as part of the normal help description. This also moves registered passes and pipelines into different sections of the help.
Example:
```
Compiler passes to run
--pass-pipeline - ...
Passes:
--affine-data-copy-generate - ...
--convert-gpu-to-spirv - ...
--workgroup-size=<long> - ...
--test-options-pass - ...
--list=<int> - ...
--string=<string> - ...
--string-list=<string> - ...
Pass Pipelines:
--test-options-pass-pipeline - ...
--list=<int> - ...
--string=<string> - ...
--string-list=<string> - ...
```
Differential Revision: https://reviews.llvm.org/D74246
Summary:
The `vector.fma` operation is portable enough across targets that we do not want
to keep it wrapped under `vector.outerproduct` and `llvm.intrin.fmuladd`.
This revision lifts the op into the vector dialect and implements the lowering to LLVM by using two patterns:
1. a pattern that lowers from n-D to (n-1)-D by unrolling when n > 2
2. a pattern that converts from 1-D to the proper LLVM representation
Reviewers: ftynse, stellaraccident, aartbik, dcaballe, jsetoain, tetuante
Reviewed By: aartbik
Subscribers: fhahn, dcaballe, merge_guards_bot, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, aartbik, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74075
Summary:
This revision exposes the portable `llvm.fma` intrinsic in LLVMOps and uses it
in lieu of `llvm.fmuladd` when lowering the `vector.outerproduct` op to LLVM.
This guarantees proper `fma` instructions will be emitted if the target ISA
supports it.
`llvm.fmuladd` does not have this guarantee in its semantics, despite evidence
that the proper x86 instructions are emitted.
For more details, see https://llvm.org/docs/LangRef.html#llvm-fmuladd-intrinsic.
Reviewers: ftynse, aartbik, dcaballe, fhahn
Reviewed By: aartbik
Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74219
The initial implementation of the fusion operation exposes a method to
fuse a consumer with its producer, when
- both the producer and consumer operate on tensors
- the producer has only a single result value
- the producer has only "parallel" iterator types
A new interface method hasTensorSemantics is added to verify that an
operation has all operands and results of type RankedTensorType.
Differential Revision: https://reviews.llvm.org/D74172
Summary: In some edge cases the default APFloat printer will generate something that we can't parse back in. In these cases, fallback to using hex instead.
Differential Revision: https://reviews.llvm.org/D74181
Technically a leak in tblgen is harmless, but this makes asan builds of
mlir very noisy. Just use a SpecificBumpPtrAllocator that knows how to
clean up after itself.
This reverts commit 64871f778d.
ASAN indicates a use-after-free in in mlir::canFuseLoops(mlir::AffineForOp, mlir::AffineForOp, unsigned int, mlir::ComputationSliceState*) lib/Transforms/Utils/LoopFusionUtils.cpp:202:41
mlir-opt needs to link against MLIRLoopAnalysis
This shouldn't be needed but MLIR "hack" for
"whole-archive" linking is not compatible with
CMake transitive dependencies management.
Differential Revision: https://reviews.llvm.org/D74097
Summary:
This revision adds basic support for emitting line table information when exporting to LLVMIR. We don't yet have a story for supporting all of the LLVM debug metadata, so this revision stubs some features(like subprograms) to enable emitting line tables.
Differential Revision: https://reviews.llvm.org/D73934
Summary:
Previously, vector.contract did not allow an empty set of
free or batch dimensions (K = 0) which defines a basic
reduction into a scalar (like a dot product). This CL
relaxes that restriction. Also adds constraints on
element type of operands and results. With tests.
Reviewers: nicolasvasilache, andydavis1, rriddle
Reviewed By: andydavis1
Subscribers: merge_guards_bot, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74014
Summary:
[mlir][VectorOps] Support vector transfer_read/write unrolling for memrefs with vector element type. When unrolling vector transfer read/write on memrefs with vector element type, the indices used to index the memref argument must be updated to reflect the unrolled operation. However, in the case of memrefs with vector element type, we need to be careful to only update the relevant memref indices.
For example, a vector transfer read with the following source/result types, memref<6x2x1xvector<2x4xf32>>, vector<2x1x2x4xf32>, should only update memref indices 1 and 2 during unrolling.
Reviewers: nicolasvasilache, aartbik
Reviewed By: nicolasvasilache, aartbik
Subscribers: lebedev.ri, Joonsoo, merge_guards_bot, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, liufengdb, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D72965
Summary:
Adds affine loop fusion transformation function to LoopFusionUtils.
Updates TestLoopFusion utility to run loop fusion transformation until a fixed point is reached.
Adds unit tests to test the transformation.
Reviewers: bondhugula, dcaballe, nicolasvasilache
Reviewed By: bondhugula, dcaballe
Subscribers: Joonsoo, merge_guards_bot, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, aartbik, liufengdb, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D73190
Summary:
Add ShapeCastOp to the vector ops dialect.
The shape_cast operation casts between an n-D source vector shape and a k-D result vector shape (the element type remains the same).
Reviewers: nicolasvasilache, aartbik
Reviewed By: nicolasvasilache
Subscribers: Joonsoo, merge_guards_bot, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, liufengdb, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D73635
Summary: Optional regions are supported in the generic op print/parse form, update the docs to match.
Differential Revision: https://reviews.llvm.org/D74061
Summary:
MLIRAnalysis depended on MLIRVectorOps
MLIRVectorOps depended on MLIRAnalysis for Loop information.
Both of these can be solved by factoring out libraries related to loop
analysis into their own library. The new MLIRLoopAnalysis might be
better off with the Loop Dialect in the future.
Reviewers: nicolasvasilache, rriddle!, mehdi_amini
Reviewed By: mehdi_amini
Subscribers: Joonsoo, vchuravy, merge_guards_bot, mgorny, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, aartbik, liufengdb, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D73655
Summary:
This breaks a cyclic library dependency where MLIRPass used the verifier
in MLIRAnalysis, but MLIRAnalysis also contained passes used for testing.
The presence of the test passes here is archaeology, predating
test/lib/Transform.
Reviewers: rriddle
Reviewed By: rriddle
Subscribers: merge_guards_bot, mgorny, mehdi_amini, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D74067
The recent refactoring of build files broke building with the MIR CUDA
integration enabled. This fixes it by adding some additional
dependencies to mlir-opt.
Differential Revision: https://reviews.llvm.org/D74041
Summary:
It is often needed to map entire ranges rather than single values. To avoid
writing the same for loop every time, I have added an overload to the map
method.
Differential Revision: https://reviews.llvm.org/D73894
This binplaces `mlir-translate`, `mlir-cuda-runner`, and `mlir-cpu-runner` when building the CMake install target.
Differential Revision: https://reviews.llvm.org/D73986
We were using normal dictionary attribute for target environment
specification. It becomes cumbersome with more and more fields.
This commit changes the modelling to a dialect-specific attribute,
where we can have control over its storage and assembly form.
Differential Revision: https://reviews.llvm.org/D73959
This is fixing a build error:
error: non-constant-expression cannot be narrowed from type 'unsigned int' to 'Region::iterator::difference_type' (aka 'int') in initializer list
Fix pr44767
Summary:
This patch is a step towards enabling BUILD_SHARED_LIBS=on, which
builds most libraries as DLLs instead of statically linked libraries.
The main effect of this is that incremental build times are greatly
reduced, since usually only one library need be relinked in response
to isolated code changes.
The bulk of this patch is fixing incorrect usage of cmake, where library
dependencies are listed under add_dependencies rather than under
target_link_libraries or under the LINK_LIBS tag. Correct usage should be
like this:
add_dependencies(MLIRfoo MLIRfooIncGen)
target_link_libraries(MLIRfoo MLIRlib1 MLIRlib2)
A separate issue is that in cmake, dependencies between static libraries
are automatically included in dependencies. In the above example, if MLIBlib1
depends on MLIRlib2, then it is sufficient to have only MLIRlib1 in the
target_link_libraries. When compiling with shared libraries, it is necessary
to have both MLIRlib1 and MLIRlib2 specified if MLIRfoo uses symbols from both.
Reviewers: mravishankar, antiagainst, nicolasvasilache, vchuravy, inouehrs, mehdi_amini, jdoerfert
Reviewed By: nicolasvasilache, mehdi_amini
Subscribers: Joonsoo, merge_guards_bot, jholewinski, mgorny, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, csigg, arpith-jacob, mgester, lucyrfox, herhut, aartbik, liufengdb, llvm-commits
Tags: #llvm
Differential Revision: https://reviews.llvm.org/D73653
This commit adds two resource limits, max_compute_workgroup_size
and max_compute_workgroup_invocations as resource limits to
the target environment. They are not used at the current moment,
but they will affect the SPIR-V CodeGen. Adding for now to have
a proper target environment modelling.
Differential Revision: https://reviews.llvm.org/D73905
This revision makes sure that errors emitted outside of testing are treated as fatal errors. This avoids the current silent failures that occur when the format is invalid.
Summary:
Currently BuildableType is assumed to be preceded by a builder. This prevents constructing types that don't have a callable 'get' method with the builder. This revision reworks the format to be like attribute builders, i.e. by accepting $_builder within the format itself.
Differential Revision: https://reviews.llvm.org/D73736
Summary: This revision add support for accepting a few type constraints, e.g. AllTypesMatch, when inferring types for operands and results. This is used to remove the c++ parsers for several additional operations.
Differential Revision: https://reviews.llvm.org/D73735
Summary:
Replace the generic zero- and one-result builders in LLVM::CallOp with a custom
builder that takes an LLVMFuncOp, which can be used to extract the result type
and create the symbol reference attribute. This is merely a convenience for
upcoming changes. The ODS-generated builders remain present.
Introduce LLVM::LLVMType::isVoidTy by analogy with the underlying LLVM type.
Differential Revision: https://reviews.llvm.org/D73895
LinalgDependenceGraph was not updated after successful producer-consumer
fusion for linalg ops. In this patch it is fixed by reconstructing
LinalgDependenceGraph on every iteration. This is very ineffective and
should be improved by updating LDGraph only when it is necessary.
The patterns for converting `std.alloc` and `std.dealoc` can be configured to
use `llvm.alloca` instead of calling `malloc` and `free`. This configuration
has been only possible through a command-line flag, despite the presence of a
(misleading) parameter in the pass constructor. Use the parameter instead and
only initalize it from the command line flags if the pass is constructed from
the mlir-opt registration.
Summary:
These hooks were originally introduced to support passes deriving the
StandardToLLVM conversion, in particular converting types from different
dialects to LLVM types in a single-step conversion. They are no longer in use
since the pass and conversion infrastructure has evolved sufficiently to make
defining new passes with exactly the same functionality simple through the use
of populate* functions, conversion targets and type converters. Remove the
hooks. Any users of this hooks can call the dialect conversion infrastructure
directly instead, which is likely to require less LoC than these hooks.
Differential Revision: https://reviews.llvm.org/D73795
Summary:
In the original design, gpu.launch required explicit capture of uses
and passing them as operands to the gpu.launch operation. This was
motivated by infrastructure restrictions rather than design. This
change lifts the requirement and removes the concept of kernel
arguments from gpu.launch. Instead, the kernel outlining
transformation now does the explicit capturing.
This is a breaking change for users of gpu.launch.
Differential Revision: https://reviews.llvm.org/D73769
Mason Remy