River Riddle [Thu, 25 Jul 2019 18:30:41 +0000 (11:30 -0700)]
Add support for an analysis mode to DialectConversion.
This mode analyzes which operations are legalizable to the given target if a conversion were to be applied, i.e. no rewrites are ever performed even on success. This mode is useful for device partitioning or other utilities that may want to analyze the effect of conversion to different targets before performing it.
The analysis method currently just fills a provided set with the operations that were found to be legalizable. This can be extended in the future to capture more information as necessary.
PiperOrigin-RevId:
259987105
Nicolas Vasilache [Thu, 25 Jul 2019 08:33:02 +0000 (01:33 -0700)]
Fix backward slice computation to iterate through known control flow
This CL fixes an oversight with dealing with loops in slicing analysis.
The forward slice computation properly propagates through loops but not the backward slice.
Add relevant unit tests.
PiperOrigin-RevId:
259903396
Alex Zinenko [Thu, 25 Jul 2019 07:40:48 +0000 (00:40 -0700)]
Move GPU dialect to {lib,include/mlir}/Dialect
Per tacit agreement, individual dialects should now live in lib/Dialect/Name
with headers in include/mlir/Dialect/Name and tests in test/Dialect/Name.
PiperOrigin-RevId:
259896851
River Riddle [Wed, 24 Jul 2019 23:41:11 +0000 (16:41 -0700)]
NFC: Use ValueOfRange instead of T in Diagnostic::appendRange.
For iterator_range, T is often the name of another iterator type and not the the value of the range.
PiperOrigin-RevId:
259843446
Alex Zinenko [Wed, 24 Jul 2019 20:13:14 +0000 (13:13 -0700)]
Move SPIRV dialect tests under test/Dialect
This was overlooked when the dialect code was moved under
{lib,include/mlir}/Dialect. NFC.
PiperOrigin-RevId:
259801927
Lei Zhang [Wed, 24 Jul 2019 19:35:24 +0000 (12:35 -0700)]
Disable auto-generated builders for spv.module
We need to ensure the block inside the region is properly terminated;
the auto-generated builders do not guarantee that.
PiperOrigin-RevId:
259793777
Nicolas Vasilache [Wed, 24 Jul 2019 17:28:04 +0000 (10:28 -0700)]
Cleanup slicing test.
Remove hardcoded SSA names and make use of CHECK-LABEL directives.
PiperOrigin-RevId:
259767803
Nicolas Vasilache [Wed, 24 Jul 2019 12:10:26 +0000 (05:10 -0700)]
Enable multi-level Linalg fusion
This CL adds support for SubViewOp in the alias analysis to permit multiple Linalg fusion passes to compose. The debugging messages are also improved for better readability. The readability benefits came in handy when tracking this issue.
A 2-level fusion test is added to capture the new behavior.
PiperOrigin-RevId:
259720246
Mahesh Ravishankar [Wed, 24 Jul 2019 05:38:23 +0000 (22:38 -0700)]
Add a utility function to populate StdOp to SPIRV Conversion Patterns
The function populateStdOpsToSPIRVPatterns appends the conversion
patterns automatically generated from StdOpsToSPIRVConversion.td to a
list of patterns
PiperOrigin-RevId:
259677890
Jacques Pienaar [Wed, 24 Jul 2019 03:09:11 +0000 (20:09 -0700)]
Update cmake files.
Use MLIR_MAIN_SRC_DIR that also works if MLIR is not checked out in projects
directory & simplify dir for EDSC test.
PiperOrigin-RevId:
259664306
MLIR Team [Tue, 23 Jul 2019 18:23:14 +0000 (11:23 -0700)]
Add sitofp to the standard dialect
Conversion from integers (window or input size, padding etc) to floating point is required to express many ML kernels, for example average pooling.
PiperOrigin-RevId:
259575284
Alex Zinenko [Tue, 23 Jul 2019 17:18:18 +0000 (10:18 -0700)]
Affine loop parallelism detection: conservatively handle unknown ops
The loop parallelism detection utility only collects the affine.load and
affine.store operations appearing inside the loop to analyze the access
patterns for the absence of dependences. However, any operation, including
unregistered operations, can appear in a body of an affine loop. If such
operation has side effects, the result of parallelism analysis is incorrect.
Conservatively assume affine loops are not parallel in presence of operations
other than affine.load, affine.store, affine.for, affine.terminator that may
have side effects.
This required to update the loop-fusion unit test that relies on parallelism
analysis and was exercising loop fusion in presence of an unregistered
operation.
PiperOrigin-RevId:
259560935
Alex Zinenko [Tue, 23 Jul 2019 17:16:52 +0000 (10:16 -0700)]
ODS: support UnitAttr in Operation definitions
A recent commit introduced UnitAttr into the ODS but did not include the
support for using UnitAttrs in operation definitions (only patterns were
supported). Extend the ODS definition of UnitAttr to be usable in operation
definition by providing a trivial builder and an accessor that returns "true"
if the unit attribute is present since the attribute presence itself has
meaning.
Additionally, test that unit attributes are effectively rewritten in patterns
in addition to the already available FileCheck tests of the generated rewriter
code.
PiperOrigin-RevId:
259560653
Alex Zinenko [Tue, 23 Jul 2019 16:26:15 +0000 (09:26 -0700)]
Introduce LLVMFuncOp
Originally, MLIR only supported functions of the built-in FunctionType. On the
conversion path to LLVM IR, we were creating MLIR functions that contained LLVM
dialect operations and used LLVM IR types for everything expect top-level
functions (e.g., a second-order function would have a FunctionType that consume
or produces a wrapped LLVM function pointer type). With MLIR functions
becoming operations, it is now possible to introduce non-built-in function
operations. This will let us use conversion patterns for function conversion,
simplify the MLIR-to-LLVM translation by removing the knowledge of the MLIR
built-in function types, and provide stronger correctness verifications (e.g.
LLVM functions only accept LLVM types).
Furthermore, we can currently construct a situation where the same function is
used with two different types: () -> () when its specified and called directly,
and !llvm<"void ()"> when it's passed somewhere on called indirectly. Having a
special function-op that is always of !llvm<"void ()"> type makes the function
model and the llvm dialect type system more consistent.
Introduce LLVMFuncOp to represent a function in the LLVM dialect. Unlike
standard FuncOp, this function has an LLVMType wrapping an LLVM IR function
type. Generalize the common behavior of function-defining operations
(functions live in a symbol table of a module, contain a single region, are
iterable as a list of blocks, and support argument attributes).
This only defines the operation. Custom syntax, conversion and translation
rules will be added in follow-ups.
The operation name mentions LLVM explicitly to avoid confusion with standard
FuncOp, especially in multiple files that use both `mlir` and `mlir::LLVM`
namespaces.
PiperOrigin-RevId:
259550940
River Riddle [Tue, 23 Jul 2019 04:43:14 +0000 (21:43 -0700)]
Allow std.constant to hold a boolean value.
This was an oversight in the original implementation, std.constant already supports IntegerAttr just not BoolAttr.
PiperOrigin-RevId:
259467710
Uday Bondhugula [Tue, 23 Jul 2019 00:41:38 +0000 (17:41 -0700)]
Introduce parser library method to parse list of region arguments
- introduce parseRegionArgumentList (similar to parseOperandList) to parse a
list of region arguments with a delimiter
- allows defining custom parse for op's with multiple/variadic number of
region arguments
- use this on the gpu.launch op (although the latter has a fixed number
of region arguments)
- add a test dialect op to test region argument list parsing (with the
no delimiter case)
Signed-off-by: Uday Bondhugula <uday@polymagelabs.com>
Closes tensorflow/mlir#40
PiperOrigin-RevId:
259442536
River Riddle [Mon, 22 Jul 2019 22:06:15 +0000 (15:06 -0700)]
Emit an error for missing '[' when parsing an AffineMapOfSSAIds.
Fixes tensorflow/mlir#51
PiperOrigin-RevId:
259415034
Lei Zhang [Mon, 22 Jul 2019 20:30:25 +0000 (13:30 -0700)]
(De)serialize composite spv.constant
This CL covers the case of composite spv.constant. We encode/decode
them into/from OpConstantComposite/OpConstantNull.
PiperOrigin-RevId:
259394700
River Riddle [Mon, 22 Jul 2019 20:22:24 +0000 (13:22 -0700)]
NFC: Update the LoopToStd conversion patterns to use RewritePattern instead of ConversionPattern.
These patterns don't require type changes so they don't need to be using ConversionPattern.
PiperOrigin-RevId:
259393151
River Riddle [Mon, 22 Jul 2019 18:28:04 +0000 (11:28 -0700)]
NFC: Update usage of multi-threading flags.
The multi-threading flags have changed as the feature is on by default and not experimental.
PiperOrigin-RevId:
259369313
Jacques Pienaar [Mon, 22 Jul 2019 17:51:40 +0000 (10:51 -0700)]
Update style/clang-format (NFC).
Update to be consistent & so that future save + clang-format workflows don't introduce extra changes.
PiperOrigin-RevId:
259361174
Lei Zhang [Mon, 22 Jul 2019 13:01:34 +0000 (06:01 -0700)]
(De)serialize float scalar spv.constant
This CL adds support for float scalar spv.constant in (de)serialization.
PiperOrigin-RevId:
259311776
Lei Zhang [Mon, 22 Jul 2019 13:01:11 +0000 (06:01 -0700)]
(De)serialize bool and integer scalar spv.constant
SPIR-V has multiple constant instructions covering different
constant types:
* `OpConstantTrue` and `OpConstantFalse` for boolean constants
* `OpConstant` for scalar constants
* `OpConstantComposite` for composite constants
* `OpConstantNull` for null constants
* ...
We model them all with a single spv.constant op for uniformity
and friendliness to transformations. This does mean that when
doing (de)serialization, we need to poke spv.constant's type
to determine which SPIR-V binary instruction to use.
This CL only covers the case of bool and integer spv.constant.
The rest will follow.
PiperOrigin-RevId:
259311698
Lei Zhang [Mon, 22 Jul 2019 13:00:47 +0000 (06:00 -0700)]
[spirv] NFC: adjust `encode*` function signatures in Serializer
* Let them return `LogicalResult` so we can chain them together
with other functions returning `LogicalResult`.
* Added "Into" as the suffix to the function name and made the
`binary` as the first parameter so that it reads more naturally.
PiperOrigin-RevId:
259311636
Alex Zinenko [Mon, 22 Jul 2019 12:37:57 +0000 (05:37 -0700)]
ODS: introduce ParamNativeOpTrait
In MLIR C++ implementation, certain OpTraits are parameterized by values or
types using nested class templates. These traits correspond to trait
identifiers of the form "OpTrait::TraitName<Parameters>::Impl". Such long
names are arguably too verbose to be spelled out entirely as NativeOpTrait in
ODS and obscure the fact of parameterization. Introduce ParamNativeOpTrait to
generate such trait identifiers from the base name and the parameters directly
in ODS. Exercise the functionality on the SingleBlockImplicitTerminator trait.
PiperOrigin-RevId:
259308796
Lei Zhang [Mon, 22 Jul 2019 12:37:12 +0000 (05:37 -0700)]
[spirv] Remove one level of indirection: processOp to processOpImpl
We already have two levels of controls in SPIRVBase.td: hasOpcode and
autogenSerialization. The former controls whether to add an entry to
the dispatch table, while the latter controls whether to autogenerate
the op's (de)serialization method specialization. This is enough for
our cases. Remove the indirection from processOp to processOpImpl
to simplify the picture.
PiperOrigin-RevId:
259308711
Nicolas Vasilache [Mon, 22 Jul 2019 11:30:50 +0000 (04:30 -0700)]
Refactor LoopParametricTiling as a test pass - NFC
This CL moves LoopParametricTiling into test/lib as a pass for purely testing purposes.
PiperOrigin-RevId:
259300264
Alex Zinenko [Mon, 22 Jul 2019 09:41:39 +0000 (02:41 -0700)]
SingleBlockImplicitTerminator: report the wrong terminator op found
In the trait verifier of SingleBlockImplicitTerminator, report the name of the
unexpected terminator op found in the end of the block in addition to the name
of the expected terminator op. This may simplify debugging, especially in
cases where the terminator is omitted for brevity and/or after a long series of
conversions.
PiperOrigin-RevId:
259287452
River Riddle [Sun, 21 Jul 2019 02:05:41 +0000 (19:05 -0700)]
Refactor region type signature conversion to be explicit via patterns.
This cl enforces that the conversion of the type signatures for regions, and thus their entry blocks, is handled via ConversionPatterns. A new hook 'applySignatureConversion' is added to the ConversionPatternRewriter to perform the desired conversion on a region. This also means that the handling of rewriting the signature of a FuncOp is moved to a pattern. A default implementation is provided via 'mlir::populateFuncOpTypeConversionPattern'. This removes the hacky implicit 'dynamically legal' status of FuncOp that was present previously, and leaves it up to the user to decide when/how to convert the signature of a function.
PiperOrigin-RevId:
259161999
Mahesh Ravishankar [Sun, 21 Jul 2019 01:11:39 +0000 (18:11 -0700)]
Add (de)serialization of EntryPointOp and ExecutionModeOp
Since the serialization of EntryPointOp contains the name of the
function as well, the function serialization emits the function name
using OpName instruction, which is used during deserialization to get
the correct function name.
PiperOrigin-RevId:
259158784
River Riddle [Sat, 20 Jul 2019 16:22:36 +0000 (09:22 -0700)]
Ensure that DenseElementAttr data is 64-bit aligned.
This allows for the raw data to be reinterpreted as the derived c++ type, e.g. ArrayRef<uint64_t>. This fixes a ubsan error for misaligned-pointer-use.
PiperOrigin-RevId:
259128031
Alex Zinenko [Sat, 20 Jul 2019 10:03:45 +0000 (03:03 -0700)]
Merge TypeUtilities library into the IR library
The TypeUtilities.{cpp,h}, currently living in {lib,include/mlir}/Support, do
not belong to the Support library. Instead, they form a separate utility
library that depends on the IR library. The operations it provides relate to
standard types (tensors, memrefs) as well as to operation manipulation, making
them a better fit for the main IR library.
PiperOrigin-RevId:
259108314
MLIR Team [Fri, 19 Jul 2019 23:10:48 +0000 (16:10 -0700)]
Fix a comment about ShapedType::getNumElements()
PiperOrigin-RevId:
259057303
Lei Zhang [Fri, 19 Jul 2019 17:32:12 +0000 (10:32 -0700)]
[spirv] Avoid printing duplicate trailing type
When printing the value attribute in spv.constant, OpAsmPrinter
already attaches a trailing type. So we don't need to duplicate
it again unless it's an array attribute, which does not have
type by default but we use it for spirv::ArrayType.
PiperOrigin-RevId:
258994197
Lei Zhang [Fri, 19 Jul 2019 16:51:08 +0000 (09:51 -0700)]
Replace bitwiseCast with llvm::bit_cast
PiperOrigin-RevId:
258986485
Lei Zhang [Fri, 19 Jul 2019 16:39:14 +0000 (09:39 -0700)]
Suppress compiler warnings regarding unused variables
Not all ops have operands or results, so it ends up there may be no
use of wordIndex or the generated op's results.
PiperOrigin-RevId:
258984485
Lei Zhang [Fri, 19 Jul 2019 16:31:55 +0000 (09:31 -0700)]
Wrap op (de)serialization methods in anonymous namespace
It's a known bug that older GCC is not happy with method specialization in
the enclosing (global) namespace:
https://gcc.gnu.org/bugzilla/show_bug.cgi?id=56480
This CL wraps the generated specialization methods in the anonymous namespace
to make sure the specialization is in the same namespace as the class.
PiperOrigin-RevId:
258983181
Jacques Pienaar [Fri, 19 Jul 2019 15:15:16 +0000 (08:15 -0700)]
Switch C++14 std::equal usage to for-loop.
Version of std::equal used required C++14, switching to for-loop for now. Just a direct change from std::equal to the equivalent using for loop.
PiperOrigin-RevId:
258970366
Alex Zinenko [Fri, 19 Jul 2019 15:13:27 +0000 (08:13 -0700)]
Move loop dialect tests into separate files - NFC
This was overlooked when moving out loop operations from Standard to a separate
dialect.
PiperOrigin-RevId:
258970115
Lei Zhang [Fri, 19 Jul 2019 14:58:04 +0000 (07:58 -0700)]
Add missing MLIRDialect dependency for MLIRDialect
Tests for the Broadcastable trait was added in a previous CL, which
makes MLIRDialect depend on MLIRDialect now.
PiperOrigin-RevId:
258967332
Mahesh Ravishankar [Fri, 19 Jul 2019 14:30:15 +0000 (07:30 -0700)]
Make SPIR-V spv.EntryPoint and spv.ExecutionMode consistent with SPIR-V spec
This CL changes the Op definition of spirv::EntryPointOp and
spirv::ExecutionModeOp to be consistent with the SPIR-V spec.
1) The EntryPointOp doesn't return a value
2) The ExecutionModeOp takes as argument, the SymbolRefAttr to refer
to the function, instead of the result of the EntryPointOp.
Following this, the spirv::EntryPointType is no longer necessary, and
is removed.
PiperOrigin-RevId:
258964027
Alex Zinenko [Fri, 19 Jul 2019 13:35:10 +0000 (06:35 -0700)]
Generalize implicit terminator into an OpTrait
Several groups of operations in different dialects (e.g. AffineForOp,
AffineIfOp; loop::ForOp, loop::IfOp) share the requirement for their regions to
contain 0 or 1 block, and for blocks to always have a specific terminator type.
Furthermore, this terminator may be omitted from the custom syntax. Generalize
this behavior into OpTrait::SingleBlockImplicitTerminator, parameterized by the
terminator operation type. This trait provides the verifier that checks the
presence of the terminator, and utility functions adding the terminator in case
of absence.
PiperOrigin-RevId:
258957180
Nicolas Vasilache [Fri, 19 Jul 2019 13:31:28 +0000 (06:31 -0700)]
Uniformize test name - NFC
PiperOrigin-RevId:
258956693
Nicolas Vasilache [Fri, 19 Jul 2019 12:02:39 +0000 (05:02 -0700)]
Refactor stripmineSink for AffineForOp - NFC
More moving less cloning.
PiperOrigin-RevId:
258947575
Nicolas Vasilache [Fri, 19 Jul 2019 11:47:27 +0000 (04:47 -0700)]
Utility function to map a loop on a parametric grid of virtual processors
This CL introduces a simple loop utility function which rewrites the bounds and step of a loop so as to become mappable on a regular grid of processors whose identifiers are given by SSA values.
A corresponding unit test is added.
For example, using CUDA terminology, and assuming a 2-d grid with processorIds = [blockIdx.x, threadIdx.x] and numProcessors = [gridDim.x, blockDim.x], the loop:
```
loop.for %i = %lb to %ub step %step {
...
}
```
is rewritten into a version resembling the following pseudo-IR:
```
loop.for %i = %lb + threadIdx.x + blockIdx.x * blockDim.x to %ub
step %gridDim.x * blockDim.x {
...
}
```
PiperOrigin-RevId:
258945942
Nicolas Vasilache [Fri, 19 Jul 2019 08:52:36 +0000 (01:52 -0700)]
Uniformize the API for the mlir::tile functions on AffineForOp and loop::ForOp
This CL adapts the recently introduced parametric tiling to have an API matching the tiling
of AffineForOp. The transformation using stripmineSink is more general and produces imperfectly nested loops.
Perfect nesting invariants of the tiled version are obtained by selectively applying hoisting of ops to isolate perfectly nested bands. Such hoisting may fail to produce a perfect loop nest in cases where ForOp transitively depend on enclosing induction variables. In such cases, the API provides a LogicalResult return but the SimpleParametricLoopTilingPass does not currently use this result.
A new unit test is added with a triangular loop for which the perfect nesting property does not hold. For this example, the old behavior was to produce IR that did not verify (some use was not dominated by its def).
PiperOrigin-RevId:
258928309
River Riddle [Fri, 19 Jul 2019 01:20:03 +0000 (18:20 -0700)]
Add support for providing a legality callback for dynamic legality in DialectConversion.
This allows for providing specific handling for dynamically legal operations/dialects without overriding the general 'isDynamicallyLegal' hook. This also means that a derived ConversionTarget class need not always be defined when some operations are dynamically legal.
Example usage:
ConversionTarget target(...);
target.addDynamicallyLegalOp<ReturnOp>([](ReturnOp op) {
return ...
};
target.addDynamicallyLegalDialect<StandardOpsDialect>([](Operation *op) {
return ...
};
PiperOrigin-RevId:
258884753
Lei Zhang [Thu, 18 Jul 2019 21:17:56 +0000 (14:17 -0700)]
[spirv] group methods better and improve comments
This CL groups (de)serialization methods logically and improves comments
at various places. It also sorted method implementations to follow the
order of their declarations. There is NFC.
PiperOrigin-RevId:
258843490
Lei Zhang [Thu, 18 Jul 2019 19:51:55 +0000 (12:51 -0700)]
Place generated StandardOps to SPIR-V patterns in anonymous namespace
This avoids polluting the mlir namespace.
PiperOrigin-RevId:
258826497
River Riddle [Thu, 18 Jul 2019 19:04:57 +0000 (12:04 -0700)]
NFC: Expose a ConversionPatternRewriter for use with ConversionPatterns.
This specific PatternRewriter will allow for exposing hooks in the future that are only useful for the conversion framework, e.g. type conversions.
PiperOrigin-RevId:
258818122
Feng Liu [Thu, 18 Jul 2019 18:25:53 +0000 (11:25 -0700)]
Add an "is_signed" attribute to the quant_ConstFakeQuant op
Some TensorFlow simulated quantize ops such as QuantizeAndDequantizeV2Op have
attribute for the sign of the quantization, so quant_ConstFakeQuant should be
able to represent it with the new attribute is added.
The method for converting these attributes to an QuantizedType is updated to
handle this new argument.
PiperOrigin-RevId:
258810290
Mehdi Amini [Thu, 18 Jul 2019 17:34:19 +0000 (10:34 -0700)]
Minor cleanup to LangRef, MLIR stands for "Multi-Level IR"
PiperOrigin-RevId:
258798577
Lei Zhang [Thu, 18 Jul 2019 16:35:36 +0000 (09:35 -0700)]
Fix script relative path after moving SPIR-V dialect
PiperOrigin-RevId:
258786729
Lei Zhang [Thu, 18 Jul 2019 16:24:41 +0000 (09:24 -0700)]
Print boolean values in ElementsAttr as "true"/"false"
We already parse boolean "true"/"false" as ElementsAttr elements.
This CL makes it round-trippable that we are printing the same way.
PiperOrigin-RevId:
258784962
Jing Pu [Thu, 18 Jul 2019 03:10:39 +0000 (20:10 -0700)]
Add UnitAttr in OpBase.td.
Note that UnitAttr cannot be used for op definition yet.
PiperOrigin-RevId:
258693338
Mahesh Ravishankar [Thu, 18 Jul 2019 01:41:28 +0000 (18:41 -0700)]
Automatically generate (de)serialization methods for SPIR-V ops
For ops in SPIR-V dialect that are a direct mirror of SPIR-V
operations, the serialization/deserialization methods can be
automatically generated from the Op specification. To enable this an
'autogenSerialization' field is added to SPV_Ops. When set to
non-zero, this will enable the automatic (de)serialization function
generation
Also adding tests that verify the spv.Load, spv.Store and spv.Variable
ops are serialized and deserialized correctly. To fully support these
tests also add serialization and deserialization of float types and
spv.ptr types
PiperOrigin-RevId:
258684764
Geoffrey Martin-Noble [Thu, 18 Jul 2019 00:06:33 +0000 (17:06 -0700)]
Add helper to get flattened tuple types
The API on TupleType::getFlattenedTypes follows our normal conventions by accepting an output parameter, but requires callers to allocate their own storage and lends itself to use in an imperative style. This makes it difficult to use in tablegen. The current solution is to define a lambda that is immediately called, but it's cleaner to extract that into a helper.
PiperOrigin-RevId:
258672046
Smit Hinsu [Wed, 17 Jul 2019 23:14:37 +0000 (16:14 -0700)]
Simplify broadcastable traits
We only verify broadcastable trait verifier and don't care about mutations so removed all CHECK statements and FileCheck invocation.
PiperOrigin-RevId:
258662882
River Riddle [Wed, 17 Jul 2019 23:05:32 +0000 (16:05 -0700)]
Add support for parsing/printing the trailing type of a dialect attribute.
This cl standardizes the printing of the type of dialect attributes to work the same as other attribute kinds. The type of dialect attributes will trail the dialect specific portion:
`#` dialect-namespace `<` attr-data `>` `:` type
The attribute parsing hooks on Dialect have been updated to take an optionally null expected type for the attribute. This matches the respective parseAttribute hooks in the OpAsmParser.
PiperOrigin-RevId:
258661298
Mehdi Amini [Wed, 17 Jul 2019 22:10:08 +0000 (15:10 -0700)]
Update Contributing.md doc to refer to the developer guide
This intends to reduce duplication and ensure that new informations added
to the developer guide (like the recent testing good practice) are not
missed by new contributors.
PiperOrigin-RevId:
258650672
Smit Hinsu [Wed, 17 Jul 2019 21:54:39 +0000 (14:54 -0700)]
Relax Broadcastable trait to only reject instances that are statically incompatible
Currently, Broadcastable trait also rejects instances when the op result has shape other than what can be statically inferred based on the operand shapes even if the result shape is compatible with the inferred broadcasted shape.
For example,
(tensor<3x2xi32>, tensor<*xi32>) -> tensor<4x3x2xi32>
(tensor<2xi32>, tensor<2xi32>) -> tensor<*xi32>
PiperOrigin-RevId:
258647493
River Riddle [Wed, 17 Jul 2019 21:45:53 +0000 (14:45 -0700)]
Refactor the conversion of block argument types in DialectConversion.
This cl begins a large refactoring over how signature types are converted in the DialectConversion infrastructure. The signatures of blocks are now converted on-demand when an operation held by that block is being converted. This allows for handling the case where a region is created as part of a pattern, something that wasn't possible previously.
This cl also generalizes the region signature conversion used by FuncOp to work on any region of any operation. This generalization allows for removing the 'apply*Conversion' functions that were specific to FuncOp/ModuleOp. The implementation currently uses a new hook on TypeConverter, 'convertRegionSignature', but this should ideally be removed in favor of using Patterns. That depends on adding support to the PatternRewriter used by ConversionPattern to allow applying signature conversions to regions, which should be coming in a followup.
PiperOrigin-RevId:
258645733
Smit Hinsu [Wed, 17 Jul 2019 21:05:19 +0000 (14:05 -0700)]
Add tests for broadcastable trait
PiperOrigin-RevId:
258637509
River Riddle [Wed, 17 Jul 2019 18:51:59 +0000 (11:51 -0700)]
Add an initial TestingGuide document to describe testing in MLIR.
As the number of contributors begins to scale, and the number of tests rise, it is important to detail the testing strategy in MLIR and best practices for writing those tests.
PiperOrigin-RevId:
258612585
River Riddle [Wed, 17 Jul 2019 16:26:57 +0000 (09:26 -0700)]
Add support for explicitly marking dialects and operations as illegal.
This explicit tag is useful is several ways:
*) This simplifies how to mark sub sections of a dialect as explicitly unsupported, e.g. my target supports all operations in the foo dialect except for these select few. This is useful for partial lowerings between dialects.
*) Partial conversions will now verify that operations that were explicitly marked as illegal must be converted. This provides some guarantee that the operations that need to be lowered by a specific pass will be.
PiperOrigin-RevId:
258582879
River Riddle [Tue, 16 Jul 2019 19:45:05 +0000 (12:45 -0700)]
Verify that ReturnOp only appears within the region of a FuncOp.
The invariants of ReturnOp are directly tied to FuncOp, making ReturnOp invalid in any other context.
PiperOrigin-RevId:
258421200
Nicolas Vasilache [Tue, 16 Jul 2019 19:20:15 +0000 (12:20 -0700)]
Move affine.for and affine.if to ODS
As the move to ODS is made, body and region names across affine and loop dialects are uniformized.
PiperOrigin-RevId:
258416590
River Riddle [Tue, 16 Jul 2019 18:57:45 +0000 (11:57 -0700)]
Refactor DialectConversion to support different conversion modes.
Users generally want several different modes of conversion. This cl refactors DialectConversion to provide two:
* Partial (applyPartialConversion)
- This mode allows for illegal operations to exist in the IR, and does not fail if an operation fails to be legalized.
* Full (applyFullConversion)
- This mode fails if any operation is not properly legalized to the conversion target. This allows for ensuring that the IR after a conversion only contains operations legal for the target.
PiperOrigin-RevId:
258412243
Jacques Pienaar [Tue, 16 Jul 2019 18:54:54 +0000 (11:54 -0700)]
Add a TypeIsPred.
Mostly one would use the type specification directly on the operand, but for
cases where the type of the operand depends on other operand types, `TypeIs`
attribute can be used to construct verification methods.
PiperOrigin-RevId:
258411758
Feng Liu [Tue, 16 Jul 2019 16:32:18 +0000 (09:32 -0700)]
Support signed and unsigned quantization types
This patch added a new argument to the fakeQuantAttrsToType utility method, so
it can be used to convert min/max to quantized type with different signed
storage types.
PiperOrigin-RevId:
258382538
Alex Zinenko [Tue, 16 Jul 2019 13:27:11 +0000 (06:27 -0700)]
Fix build by making LoopOps depend on StandardOps
LoopOps needs the definition ConstantIndexOp in the verifier of loop::ForOp.
PiperOrigin-RevId:
258355329
Stephan Herhut [Tue, 16 Jul 2019 12:28:56 +0000 (05:28 -0700)]
Move shared cpu runner library to Support/JitRunner.
PiperOrigin-RevId:
258347825
Lei Zhang [Tue, 16 Jul 2019 12:06:57 +0000 (05:06 -0700)]
NFC: Move SPIR-V dialect to Dialect/ subdirectory
PiperOrigin-RevId:
258345603
Lei Zhang [Tue, 16 Jul 2019 11:33:54 +0000 (04:33 -0700)]
Better support for attribute wrapper classes when getting def name
Unless we explicitly name a template instantiation in .td file, its def
name will be "anonymous_<number>". We typically give base-level Attr
template instantiation a name by writing `def AnAttr : Attr<...>`. But
when `AnAttr` is further wrapped in classes like OptionalAttr, the name
is lost unless explicitly def'ed again. These implicit-named template
instantiation is fairly common when writing op definitions. Those wrapper
classes are just essentially attaching more information to the attribute.
Without a proper way to trace back to the original attribute def name
can cause problems for consumers wanting to handle attributes according
to their types.
Previously we handled OptionalAttr and DefaultValuedAttr specifically,
but Confined was not supported. And they can compose together to have
Confined<OptionalAttr<...>, [...]>. So this CL moves the baseAttr field
to main Attr class (like isOptional) and set it only on the innermost
wrapper class.
PiperOrigin-RevId:
258341646
Nicolas Vasilache [Tue, 16 Jul 2019 08:46:23 +0000 (01:46 -0700)]
Replace linalg.for by loop.for
With the introduction of the Loop dialect, uses of the `linalg.for` operation can now be subsumed 1-to-1 by `loop.for`.
This CL performs the replacement and tests are updated accordingly.
PiperOrigin-RevId:
258322565
Alex Zinenko [Tue, 16 Jul 2019 08:27:08 +0000 (01:27 -0700)]
Forward-declare LogicalResult as struct rather than class
Windows builds are broken by a class/struct mismatch between a
forward-declaration of LogicalResult and its definition.
PiperOrigin-RevId:
258320420
River Riddle [Mon, 15 Jul 2019 19:52:44 +0000 (12:52 -0700)]
Remove lowerAffineConstructs and lowerControlFlow in favor of providing patterns.
These methods don't compose well with the rest of conversion framework, and create artificial breaks in conversion. Replace these methods with two(populateAffineToStdConversionPatterns and populateLoopToStdConversionPatterns respectively) that populate a list of patterns to perform the same behavior.
PiperOrigin-RevId:
258219277
River Riddle [Mon, 15 Jul 2019 16:52:52 +0000 (09:52 -0700)]
Update 'applyPatternsGreedily' to work on the regions of any operations.
'applyPatternsGreedily' is a useful utility outside of just function regions.
PiperOrigin-RevId:
258182937
River Riddle [Mon, 15 Jul 2019 15:40:11 +0000 (08:40 -0700)]
Refactor the traversal of operations to Convert in DialectConversion.
This cl changes the way that operations/blocks to convert are collected/traversed so that parent region operations can be legalized before their bodies. Most RewritePatterns for region operations assume that the entry arguments to each region are yet to be converted. Given that the bodies are currently converted first, this makes it difficult to fit these patterns into the same run as one converting types.
The operations/blocks to convert are now collected before any legalization has run, which simplifies the conversion logic itself, as legalization may insert new operations, move blocks, etc.
PiperOrigin-RevId:
258170158
Andy Davis [Mon, 15 Jul 2019 15:29:47 +0000 (08:29 -0700)]
Fix opt build (unused variable in Linalg).
PiperOrigin-RevId:
258168108
Alex Zinenko [Mon, 15 Jul 2019 14:50:11 +0000 (07:50 -0700)]
Linalg Utils: use Doxygen comments where appropriate
PiperOrigin-RevId:
258160982
Alex Zinenko [Mon, 15 Jul 2019 14:42:37 +0000 (07:42 -0700)]
Extend linalg transformations to allow value operands that are not views
This CL extends the linalg ops that can be tiled and fused to operations that take either views, scalar or vector operands.
PiperOrigin-RevId:
258159734
River Riddle [Mon, 15 Jul 2019 14:35:00 +0000 (07:35 -0700)]
Refactor LowerAffine to use OpRewritePattern instead of ConversionPattern.
ConversionPattern should ideally only be used when the types of the operands are changing, which in this case they aren't. Using OpRewritePattern also lends to much simpler code.
PiperOrigin-RevId:
258158474
Alex Zinenko [Mon, 15 Jul 2019 13:55:53 +0000 (06:55 -0700)]
LLVMDialect: still depend on standard types.
The dependecy on standard type is mandated by attribute verifiers on some
operations (e.g., llvm.constant) that use values of those types.
PiperOrigin-RevId:
258153134
MLIR Team [Mon, 15 Jul 2019 13:44:38 +0000 (06:44 -0700)]
Automated rollback of changelist
258149291.
PiperOrigin-RevId:
258151618
Nicolas Vasilache [Mon, 15 Jul 2019 13:44:22 +0000 (06:44 -0700)]
Start moving AffineOps to ODS
This CL starts with the AffineTerminatorOp to ODS.
PiperOrigin-RevId:
258151579
Alex Zinenko [Mon, 15 Jul 2019 13:40:07 +0000 (06:40 -0700)]
Introduce loop coalescing utility and a simple pass
Multiple (perfectly) nested loops with independent bounds can be combined into
a single loop and than subdivided into blocks of arbitrary size for load
balancing or more efficient parallelism exploitation. However, MLIR wants to
preserve the multi-dimensional multi-loop structure at higher levels of
abstraction. Introduce a transformation that coalesces nested loops with
independent bounds so that they can be further subdivided by tiling.
PiperOrigin-RevId:
258151016
Nicolas Vasilache [Mon, 15 Jul 2019 13:27:21 +0000 (06:27 -0700)]
Extend linalg transformations to allow value operands that are not views
This CL extends the linalg ops that can be tiled and fused to operations that take either views, scalar or vector operands.
PiperOrigin-RevId:
258149291
Alex Zinenko [Mon, 15 Jul 2019 13:19:39 +0000 (06:19 -0700)]
Decouple LLVM dialect from Standard dialect
Due to the absence of ODS support for enum attributes, the implementation of
the LLVM dialect `icmp` operation was reusing the comparison predicate from the
Standard dialect, creating an avoidable library dependency. With ODS support
and ICmpPredicate attribute recently introduced, the dependency is no longer
justified. Update the Standard to LLVM convresion to also convert the
CmpIPredicate into LLVM::ICmpPredicate and remove the unnecessary includes.
Note that the MLIRLLVMIR library did not explicitly depend on MLIRStandardOps,
requiring dependees of MLIRLLVMIR to also depend on MLIRStandardOps, which
should no longer be the case.
PiperOrigin-RevId:
258148456
Alex Zinenko [Mon, 15 Jul 2019 12:45:08 +0000 (05:45 -0700)]
LLVM Dialect: define ICmpPredicate in ODS
Use the recently introduced enum-gen functionality to define the predicate
attribute of the ICmp LLVM dialect operation directly in ODS. This removes the
need for manually-coded string-to-integer conversion functions and contributes
to the overall homogenization of the operation definitions.
PiperOrigin-RevId:
258143923
Nicolas Vasilache [Mon, 15 Jul 2019 09:50:09 +0000 (02:50 -0700)]
Extract std.for std.if and std.terminator in their own dialect
These ops should not belong to the std dialect.
This CL extracts them in their own dialect and updates the corresponding conversions and tests.
PiperOrigin-RevId:
258123853
River Riddle [Sun, 14 Jul 2019 20:29:46 +0000 (13:29 -0700)]
Fix a bug in DialectConversion when using RewritePattern.
When using a RewritePattern and replacing an operation with an existing value, that value may have already been replaced by something else. This cl ensures that only the final value is used when applying rewrites.
PiperOrigin-RevId:
258058488
Masaki Kozuki [Sat, 13 Jul 2019 12:55:25 +0000 (05:55 -0700)]
Fix typos
Closes tensorflow/mlir#45
PiperOrigin-RevId:
257948893
River Riddle [Sat, 13 Jul 2019 01:11:40 +0000 (18:11 -0700)]
NFC: Add header blocks to DialectConversion.h to improve readability.
PiperOrigin-RevId:
257903383
River Riddle [Fri, 12 Jul 2019 22:49:38 +0000 (15:49 -0700)]
NFC: Don't print the location of a diagnostic if it is unknown.
Printing 'loc(unknown)' just clutters the output with unhelpful information.
PiperOrigin-RevId:
257883717
Mahesh Ravishankar [Fri, 12 Jul 2019 21:32:20 +0000 (14:32 -0700)]
Add serialization and deserialization of FuncOps. To support this the
following SPIRV Instructions serializaiton/deserialization are added
as well
OpFunction
OpFunctionParameter
OpFunctionEnd
OpReturn
PiperOrigin-RevId:
257869806
Lei Zhang [Fri, 12 Jul 2019 20:53:57 +0000 (13:53 -0700)]
[spirv] Various small code improvements
* Changed SPIR-V types to all use unsigned for member type count and index
* Used the same method name for getting element type and count
* Improved spv.CompositeExtract verification a bit
PiperOrigin-RevId:
257862580
River Riddle [Fri, 12 Jul 2019 19:11:24 +0000 (12:11 -0700)]
Update the PatternRewriter constructor to take a context instead of a region.
This will allow for cleanly using a rewriter for multiple different regions.
PiperOrigin-RevId:
257845371
River Riddle [Fri, 12 Jul 2019 18:20:09 +0000 (11:20 -0700)]
Change the IR printing pass instrumentation to ignore the verifier passes on non-failure.
The verifier passes are NO-OP and are only useful to print after in the case of failure. This removes a lot of unnecessary clutter when printing after/before all passes.
PiperOrigin-RevId:
257836310
River Riddle [Fri, 12 Jul 2019 17:43:11 +0000 (10:43 -0700)]
Remove the 'region' field from OpBuilder.
This field wasn't updated as the insertion point changed, making it potentially dangerous given the multi-level of MLIR(e.g. 'createBlock' would always insert the new block in 'region'). This also allows for building an OpBuilder with just a context.
PiperOrigin-RevId:
257829135