Commit Graph

215 Commits

Author SHA1 Message Date
River Riddle cb1777127c [mlir] Remove successor operands from the Operation class
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
2020-03-05 12:53:02 -08:00
River Riddle 621d7cca37 [mlir] Add a new BranchOpInterface to allow for opaquely interfacing with branching terminator operations.
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
2020-03-05 12:50:35 -08:00
River Riddle c0fd5e657e [mlir] Add traits for verifying the number of successors and providing relevant accessors.
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
2020-03-05 12:49:59 -08:00
Lei Zhang 9600b55ac8 [mlir][spirv] Support integer signedness
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
2020-03-04 15:14:11 -05:00
Lei Zhang 440ef33073 [mlir][spirv] Add get() method to TargetEnvAttr taking raw values
Also make the getResourceLimits() method more explicit about its return type.

Differential Revision: https://reviews.llvm.org/D75484
2020-03-04 14:01:26 -05:00
Stephen Neuendorffer 798e661567 Revert "[MLIR] Move from using target_link_libraries to LINK_LIBS for llvm libraries."
This reverts commit 7a6c689771.
This breaks the build with cmake 3.13.4, but succeeds with cmake 3.15.3
2020-02-29 11:52:08 -08:00
Stephen Neuendorffer d675df0379 Revert "[MLIR] Move from add_dependencies() to DEPENDS"
This reverts commit 31e07d716a.
2020-02-29 11:52:08 -08:00
Stephen Neuendorffer 31e07d716a [MLIR] Move from add_dependencies() to DEPENDS
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
2020-02-29 10:47:27 -08:00
Stephen Neuendorffer 7a6c689771 [MLIR] Move from using target_link_libraries to LINK_LIBS for llvm libraries.
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
2020-02-29 10:47:26 -08:00
Stephen Neuendorffer dc1056a3f1 Revert "[MLIR] Move from using target_link_libraries to LINK_LIBS for llvm libraries."
This reverts commit 2f265e3528.
2020-02-28 14:13:30 -08:00
Stephen Neuendorffer 67f2a43cf8 Revert "[MLIR] Move from add_dependencies() to DEPENDS"
This reverts commit 8a2b86b2c2.
2020-02-28 12:17:40 -08:00
Stephen Neuendorffer 8a2b86b2c2 [MLIR] Move from add_dependencies() to DEPENDS
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
2020-02-28 11:35:18 -08:00
Stephen Neuendorffer 2f265e3528 [MLIR] Move from using target_link_libraries to LINK_LIBS for llvm libraries.
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
2020-02-28 11:35:17 -08:00
Stephen Neuendorffer 5869552821 [MLIR] Refactor handling of dialect libraries
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
2020-02-28 11:35:16 -08:00
Lei Zhang 63779fb462 [mlir][spirv] Refactoring to avoid calling the same function twice 2020-02-26 15:36:54 -05:00
Lei Zhang 5bc6ff6455 [mlir][spirv] Add some folders for spv.LogicalAnd/spv.LogicalOr
This commit handles folding spv.LogicalAnd/spv.LogicalOr when
one of the operands is constant true/false.

Differential Revision: https://reviews.llvm.org/D75195
2020-02-26 15:13:37 -05:00
Lei Zhang 1e9321e97a [mlir][spirv] NFC: move folders and canonicalizers in a separate file
This gives us better file organization and faster compilation time
by avoid having a gigantic SPIRVOps.cpp file.
2020-02-26 12:41:14 -05:00
River Riddle 9eb436feaa [mlir][DeclarativeParser] Add support for formatting the successors of an operation.
This revision add support for formatting successor variables in a similar way to operands, attributes, etc.

Differential Revision: https://reviews.llvm.org/D74789
2020-02-21 15:15:32 -08:00
River Riddle b1de971ba8 [mlir][ODS] Add support for specifying the successors of an operation.
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
2020-02-21 15:15:32 -08:00
Hanhan Wang 29ad9d6b26 [mlir][spirv] Add lowering for load/store zero-rank memref from std to SPIR-V.
Differential Revision: https://reviews.llvm.org/D74874
2020-02-21 14:41:12 -05:00
Lei Zhang 35b685270b [mlir] Add a signedness semantics bit to IntegerType
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
2020-02-21 09:16:54 -05:00
River Riddle 70d8fec7c9 [mlir] Refactor the structure of the 'verifyConstructionInvariants' methods.
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
2020-02-20 10:37:52 -08:00
River Riddle 0d7ff220ed [mlir] Refactor TypeConverter to add conversions without inheritance
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
2020-02-18 16:17:48 -08:00
River Riddle 5756bc4382 [mlir][DeclarativeParser] Add support for formatting enum attributes in the string form.
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
2020-02-13 17:11:48 -08:00
Lei Zhang d3e7816d85 [mlir][spirv] Introduce spv.func
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
2020-02-12 07:46:43 -05:00
Mehdi Amini c64770506b Remove static registration for dialects, and the "alwayslink" hack for passes
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
2020-02-12 09:13:02 +00:00
Lei Zhang 50aeeed8a2 [mlir][spirv] Use spv.entry_point_abi in GPU to SPIR-V conversions
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
2020-02-10 16:24:48 -05:00
River Riddle 2f94ce0dcf [mlir][DeclarativeParser] Move several missed parsers over to the declarative form.
Differential Revision: https://reviews.llvm.org/D74283
2020-02-08 15:47:55 -08:00
Lei Zhang 13b197c7d1 [mlir][spirv] Add dialect-specific attribute for target environment
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
2020-02-04 21:33:13 -05:00
Stephen Neuendorffer d7cbef2714 [MLIR] Fixes for shared library dependencies.
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
2020-02-04 08:56:37 -08:00
Lei Zhang aad352f77c [mlir][spirv] Wrap debug-only method in #ifndef NDEBUG 2020-02-04 08:57:29 -05:00
Lei Zhang 399887c9e4 [mlir][spirv] Add resource limits into target environment
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
2020-02-04 08:35:19 -05:00
River Riddle 7ef37a5f99 [mlir] Initial support for type constraints in the declarative assembly format
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
2020-02-03 21:55:09 -08:00
Lei Zhang df71000d7d [mlir][spirv] Convert linalg.generic for reduction to SPIR-V ops
This commit adds a pattern to lower linalg.generic for reduction
to spv.GroupNonUniform* ops. Right now this only supports integer
reduction on 1-D input memref. Shader entry point ABI is queried
to make sure that the input memref's shape matches the local
workgroup's invocation configuration. This makes sure that the
workload fits in one local workgroup so that we can leverage
SPIR-V group non-uniform operations.

linglg.generic is a structured op that preserves the right level
of information. It is easier to recognize reduction at this level
than performing analysis on loops.

This commit also exposes `getElementPtr` in SPIRVLowering.h given
that it's a generally useful utility function.

Differential Revision: https://reviews.llvm.org/D73437
2020-01-31 09:37:04 -05:00
River Riddle 389b126210 [mlir][NFC] Update several SPIRV operations to use declarative parsers.
Differential Revision: https://reviews.llvm.org/D73504
2020-01-30 11:43:41 -08:00
Denis Khalikov 731b140a52 [mlir][spirv] Add GroupNonUniform arithmetic operations.
Add GroupNonUniform arithmetic operations: FAdd, FMul, IMul.
Unify parser, printer, verifier for GroupNonUniform arithmetic
operations.

Differential Revision: https://reviews.llvm.org/D73491
2020-01-28 10:21:56 -05:00
River Riddle aff4ed7326 [mlir][NFC] Update Operation::getResultTypes to use ArrayRef<Type> instead of iterator_range.
Summary: The new internal representation of operation results now allows for accessing the result types to be more efficient. Changing the API to ArrayRef is more efficient and removes the need to explicitly materialize vectors in several places.

Differential Revision: https://reviews.llvm.org/D73429
2020-01-27 19:57:48 -08:00
Lei Zhang 29e411b3d6 [mlir] Expose getNearestSymbolTable as SymbolTable class method
This is a generally useful utility function for interacting with
symbol tables.

Differential Revision: https://reviews.llvm.org/D73433
2020-01-26 17:35:26 -05:00
Lei Zhang 8d6884a15e [mlir][spirv] Create builtin variable in nearest symbol table
This commit changes the logic of `getBuiltinVariableValue` to get
or create the builtin variable in the nearest symbol table. This
will allow us to use this function in other partial conversion
cases where we haven't created the spv.module yet.

Differential Revision: https://reviews.llvm.org/D73416
2020-01-26 11:00:49 -05:00
Lei Zhang 09f9deaff2 [mlir][spirv] NFC: simplify load/store builder call sites
This commit introduces default values for load/store builders to
simplify builder call sites.

Differential Revision: https://reviews.llvm.org/D73419
2020-01-26 10:45:42 -05:00
Lei Zhang 91d6655a29 [mlir][spirv] NFC: expose builtin func op conversion pattern
This commit exposes the func op conversion pattern via a new
`populateBuiltinFuncToSPIRVPatterns` function from the standard
to SPIR-V conversion passs. This is structurally better given
that func op belongs to the builtin dialect. More importantly,
this makes the pattern reusable to other dialect to SPIR-V
dialect conversion as other dialect can well adopt builtin
func op instead of having its own. Besides, it's very common
to use func ops as test wrappers in lit tests, so test passes
will need to handle func ops too.

Differential Revision: https://reviews.llvm.org/D73421
2020-01-26 10:42:06 -05:00
Lei Zhang 60d541e1b9 [mlir][spirv] Relax verification to allow flexible placement
Thus far certain SPIR-V ops have been required to be in spv.module.
While this provides strong verification to catch unexpected errors,
it's quite rigid and makes progressive lowering difficult. Sometimes
we would like to partially lower ops from other dialects, which may
involve creating ops like global variables that should be placed in
other module-like ops. So this commit relaxes the requirement of
such SPIR-V ops' scope to module-like ops. Similarly for function-
like ops.

Differential Revision: https://reviews.llvm.org/D73415
2020-01-26 10:39:45 -05:00
Lei Zhang ae21e37eb4 [mlir][spirv] Add spv.GroupNonUniformElect and spv.GroupNonUniformIAdd
Differential Revision: https://reviews.llvm.org/D73349
2020-01-26 10:20:40 -05:00
Mehdi Amini 308571074c Mass update the MLIR license header to mention "Part of the LLVM project"
This is an artifact from merging MLIR into LLVM, the file headers are
now aligned with the rest of the project.
2020-01-26 03:58:30 +00:00
Denis Khalikov 4460cb5bcd [mlir][spirv] Add lowering for composite std.constant.
Add lowering for constant operation with ranked tensor type to
spv.constant with spv.array type.

Differential Revision: https://reviews.llvm.org/D73022
2020-01-22 08:25:00 -05:00
Denis Khalikov 3023352a7d [mlir][spirv] Simplify scalar type size calculation.
Simplify scalar type size calculation and reject boolean memrefs.

Differential Revision: https://reviews.llvm.org/D72999
2020-01-21 12:15:37 -05:00
Kazuaki Ishizaki fc817b09e2 [mlir] NFC: Fix trivial typos in comments
Differential Revision: https://reviews.llvm.org/D73012
2020-01-20 03:17:03 +00:00
Lei Zhang 927f8f40a4 [mlir][spirv] Explicitly construct ArrayRef from static array
Again for pleasing GCC 5.
2020-01-17 14:53:51 -05:00
Lei Zhang 8bcf976841 [mlir][spirv] Add `const` qualifier for static arrays
This makes the local variable `implies` to have the correct
type to satisfy ArrayRef's constructor:

  /*implicit*/ constexpr ArrayRef(const T (&Arr)[N])

Hopefully this should please GCC 5.

Differential Revision: https://reviews.llvm.org/D72924
2020-01-17 11:45:41 -05:00
Lei Zhang 267483ac70 [mlir][spirv] Support implied extensions and capabilities
In SPIR-V, when a new version is introduced, it is possible some
existing extensions will be incorporated into it so that it becomes
implicitly declared if targeting the new version. This affects
conversion target specification because we need to take this into
account when allowing what extensions to use.

For a capability, it may also implies some other capabilities,
for example, the `Shader` capability implies `Matrix` the capability.
This should also be taken into consideration when preparing the
conversion target: when we specify an capability is allowed, all
its recursively implied capabilities are also allowed.

This commit adds utility functions to query implied extensions for
a given version and implied capabilities for a given capability
and updated SPIRVConversionTarget to use them.

This commit also fixes a bug in availability spec. When a symbol
(op or enum case) can be enabled by an extension, we should drop
it's minimal version requirement. Being enabled by an extension
naturally means the symbol can be used by *any* SPIR-V version
as long as the extension is supported. The grammar still encodes
the 'version' field for such cases, but it should be interpreted
as a different way: rather than meaning a minimal version
requirement, it says the symbol becomes core at that specific
version.

Differential Revision: https://reviews.llvm.org/D72765
2020-01-17 08:01:57 -05:00