platform/upstream/llvm.git
5 years ago[spirv] Add binary arithmetic operations.
Denis Khalikov [Tue, 30 Jul 2019 16:42:33 +0000 (09:42 -0700)]
[spirv] Add binary arithmetic operations.

Add binary operations such as: OpIAdd, OpFAdd, OpISub, OpFSub, OpIMul,
OpFDiv, OpFRem, OpFMod.

Closes tensorflow/mlir#54

PiperOrigin-RevId: 260734166

5 years agoRemove dead code.
Jacques Pienaar [Mon, 29 Jul 2019 21:52:51 +0000 (14:52 -0700)]
Remove dead code.

PiperOrigin-RevId: 260585594

5 years agoRewriterGen: properly handle zero-result ops
Alex Zinenko [Mon, 29 Jul 2019 21:12:15 +0000 (14:12 -0700)]
RewriterGen: properly handle zero-result ops

RewriterGen was emitting invalid C++ code if the pattern required to create a
zero-result operation due to the absence of a special case that would avoid
generating a spurious comma.  Handle this case.  Also add rewriter tests for
zero-argument operations.

PiperOrigin-RevId: 260576998

5 years agoFix SingleBlockImplicitTerminator traits to catch empty blocks
Mehdi Amini [Mon, 29 Jul 2019 18:21:55 +0000 (11:21 -0700)]
Fix SingleBlockImplicitTerminator traits to catch empty blocks

The code was written with the assumption that on failure an error would be
issued by another verifier. However verification is stopping on the first
failure which lead to an empty output. Instead we make sure an error is
displayed.
Also add tests in the test dialect for this trait.

PiperOrigin-RevId: 260541290

5 years agoSimplify ODS for loop.if and loop.for traits (NFC)
Mehdi Amini [Mon, 29 Jul 2019 18:14:25 +0000 (11:14 -0700)]
Simplify ODS for loop.if and loop.for traits (NFC)

There is a wrapper for SingleBlockImplicitTerminator in ODS, this is nicer to
read than using `NativeOpTrait`.

PiperOrigin-RevId: 260539473

5 years agoEnable (de)serialization support for spirv::AccessChainOp
Mahesh Ravishankar [Mon, 29 Jul 2019 17:45:17 +0000 (10:45 -0700)]
Enable (de)serialization support for spirv::AccessChainOp

Automatic generation of spirv::AccessChainOp (de)serialization needs
the (de)serialization emitters to handle argument specified as
Variadic<...>. To handle this correctly, this argument can only be
the last entry in the arguments list.
Add a test to (de)serialize spirv::AccessChainOp

PiperOrigin-RevId: 260532598

5 years agoAdd a `HasParent` operation trait to enforce a specific parent on an operation (NFC)
Mehdi Amini [Mon, 29 Jul 2019 17:45:17 +0000 (10:45 -0700)]
Add a `HasParent` operation trait to enforce a specific parent on an operation (NFC)

PiperOrigin-RevId: 260532592

5 years ago[mlir-tblgen] Emit forward declarations for all the classes before the definitions
Mehdi Amini [Mon, 29 Jul 2019 17:44:33 +0000 (10:44 -0700)]
[mlir-tblgen] Emit forward declarations for all the classes before the definitions

This allows classes to refer to each other in the ODS file, for instance for traits.

PiperOrigin-RevId: 260532419

5 years agoInitialize union to avoid -Wmissing-field-initializers warning.
Jacques Pienaar [Sat, 27 Jul 2019 18:46:22 +0000 (11:46 -0700)]
Initialize union to avoid -Wmissing-field-initializers warning.

Reported by clang-6.

PiperOrigin-RevId: 260311814

5 years agoVerify that affine.load/store/dma_start/dma_wait operands are valid dimension or...
Andy Davis [Fri, 26 Jul 2019 20:00:01 +0000 (13:00 -0700)]
Verify that affine.load/store/dma_start/dma_wait operands are valid dimension or symbol identifiers.

PiperOrigin-RevId: 260197567

5 years agoAutomated rollback of commit 3708f53219aa2b201e82e7172c5064c1eb9d483b
Nicolas Vasilache [Fri, 26 Jul 2019 13:49:41 +0000 (06:49 -0700)]
Automated rollback of commit 3708f53219aa2b201e82e7172c5064c1eb9d483b

PiperOrigin-RevId: 260136255

5 years agoAdd sgemm specializations - NFC
Nicolas Vasilache [Fri, 26 Jul 2019 12:40:58 +0000 (05:40 -0700)]
Add sgemm specializations - NFC

This CL adds a few specializations for sgemm.
A minor change to alpha is made in cblas_interface.cpp to be compatible with actual BLAS calls.

For now this is for internal testing purposes only.

PiperOrigin-RevId: 260129027

5 years agoSupport referencing a single value generated by a matched multi-result op
Lei Zhang [Fri, 26 Jul 2019 11:31:15 +0000 (04:31 -0700)]
Support referencing a single value generated by a matched multi-result op

It's quite common that we want to put further constraints on the matched
multi-result op's specific results. This CL enables referencing symbols
bound to source op with the `__N` syntax.

PiperOrigin-RevId: 260122401

5 years agoFix backward slice corner case
Nicolas Vasilache [Fri, 26 Jul 2019 10:48:51 +0000 (03:48 -0700)]
Fix backward slice corner case

In the backward slice computation, BlockArgument coming from function arguments represent a natural boundary for the traversal and should not trigger llvm_unreachable.
This CL also improves the error message and adds a relevant test.

PiperOrigin-RevId: 260118630

5 years agoFix linalg_matmul_impl interfacing with sgemm
Nicolas Vasilache [Fri, 26 Jul 2019 10:33:53 +0000 (03:33 -0700)]
Fix linalg_matmul_impl interfacing with sgemm

This CL provides a fix that makes linal_matmul_impl compliant with the BLAS interface. Before this CL it would compute either C += A * B when called with cblas.cpp:cblas_sgemm implementation and C = A * B with other implementations.

PiperOrigin-RevId: 260117367

5 years agoUse "standard" load and stores in LowerVectorTransfers
Nicolas Vasilache [Fri, 26 Jul 2019 09:33:58 +0000 (02:33 -0700)]
Use "standard" load and stores in LowerVectorTransfers

Clipping creates non-affine memory accesses, use std_load and std_store instead of affine_load and affine_store.
In the future we may also want a fill with the neutral element rather than clip, this would make the accesses affine if we wanted more analyses and transformations to happen post lowering to pointwise copies.

PiperOrigin-RevId: 260110503

5 years agoAdd API for bulk-setting attributes in an OperationState
Krzysztof Drewniak [Thu, 25 Jul 2019 23:13:18 +0000 (16:13 -0700)]
Add API for bulk-setting attributes in an OperationState

This API is needed for the TFLite flatbuffer importer as the importer constructs
arrays of NamedAttributes based on the flatbuffer that need to be added to an
operation.

PiperOrigin-RevId: 260041403

5 years agoAutomated rollback of commit fc194a4f22fe53f46729821d9c4a993fe200facf
Mehdi Amini [Thu, 25 Jul 2019 22:50:54 +0000 (15:50 -0700)]
Automated rollback of commit fc194a4f22fe53f46729821d9c4a993fe200facf

PiperOrigin-RevId: 260037115

5 years ago[spirv] Add AccessChainOp operation.
Denis Khalikov [Thu, 25 Jul 2019 22:42:41 +0000 (15:42 -0700)]
[spirv] Add AccessChainOp operation.

AccessChainOp creates a pointer into a composite object that can be used with
OpLoad and OpStore.

Closes tensorflow/mlir#52

PiperOrigin-RevId: 260035676

5 years agoUse perfect forwarding for OpBuilder::create args (NFC)
Mehdi Amini [Thu, 25 Jul 2019 22:27:31 +0000 (15:27 -0700)]
Use perfect forwarding for OpBuilder::create args (NFC)

This looks like an oversight, and it can be useful for building using
non-copyable types.

PiperOrigin-RevId: 260032944

5 years agoGenericize function-like printer and parser. NFC
Alex Zinenko [Thu, 25 Jul 2019 21:26:41 +0000 (14:26 -0700)]
Genericize function-like printer and parser. NFC

Function-like operations are likely to have similar custom syntax, in
particular they all need to print function signature with argument attributes.

Transform function printer and parser so that they can be applied to any
operation with the FunctionLike trait.  Move them to the trait itself.  To
avoid large member functions in the class template, define a concrete base
class for the trait and implement common functionality in it.  This allows
printer and parser to be implemented in a source file without templating.

PiperOrigin-RevId: 260020893

5 years agoAdd support for hexadecimal float literals
Alex Zinenko [Thu, 25 Jul 2019 21:15:33 +0000 (14:15 -0700)]
Add support for hexadecimal float literals

MLIR does not have support for parsing special floating point values such as
infinities and NaNs.  If programmatically constructed, these values are printed
as NaN and (+-)Inf and cannot be parsed back.  Add parser support for
hexadecimal literals in float attributes, following LLVM IR.  The literal
corresponds to the in-memory representation of the floating point value.
IEEE 754 defines a range of possible values for NaNs, storing the bitwise
representation allows MLIR to properly roundtrip NaNs with different bit values
of significands.

PiperOrigin-RevId: 260018802

5 years agoAdd support for an analysis mode to DialectConversion.
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

5 years agoFix backward slice computation to iterate through known control flow
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

5 years agoMove GPU dialect to {lib,include/mlir}/Dialect
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

5 years agoNFC: Use ValueOfRange instead of T in Diagnostic::appendRange.
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

5 years agoMove SPIRV dialect tests under test/Dialect
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

5 years agoDisable auto-generated builders for spv.module
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

5 years agoCleanup slicing test.
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

5 years agoEnable multi-level Linalg fusion
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

5 years agoAdd a utility function to populate StdOp to SPIRV Conversion Patterns
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

5 years agoUpdate cmake files.
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

5 years agoAdd sitofp to the standard dialect
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

5 years agoAffine loop parallelism detection: conservatively handle unknown ops
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

5 years agoODS: support UnitAttr in Operation definitions
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

5 years agoIntroduce LLVMFuncOp
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

5 years agoAllow std.constant to hold a boolean value.
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

5 years agoIntroduce parser library method to parse list of region arguments
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

5 years agoEmit an error for missing '[' when parsing an AffineMapOfSSAIds.
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

5 years ago(De)serialize composite spv.constant
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

5 years agoNFC: Update the LoopToStd conversion patterns to use RewritePattern instead of Conver...
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

5 years agoNFC: Update usage of multi-threading flags.
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

5 years agoUpdate style/clang-format (NFC).
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

5 years ago(De)serialize float scalar spv.constant
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

5 years ago(De)serialize bool and integer scalar spv.constant
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

5 years ago[spirv] NFC: adjust `encode*` function signatures in Serializer
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

5 years agoODS: introduce ParamNativeOpTrait
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

5 years ago[spirv] Remove one level of indirection: processOp to processOpImpl
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

5 years agoRefactor LoopParametricTiling as a test pass - NFC
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

5 years agoSingleBlockImplicitTerminator: report the wrong terminator op found
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

5 years agoRefactor region type signature conversion to be explicit via patterns.
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

5 years agoAdd (de)serialization of EntryPointOp and ExecutionModeOp
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

5 years agoEnsure that DenseElementAttr data is 64-bit aligned.
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

5 years agoMerge TypeUtilities library into the IR library
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

5 years agoFix a comment about ShapedType::getNumElements()
MLIR Team [Fri, 19 Jul 2019 23:10:48 +0000 (16:10 -0700)]
Fix a comment about ShapedType::getNumElements()

PiperOrigin-RevId: 259057303

5 years ago[spirv] Avoid printing duplicate trailing type
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

5 years agoReplace bitwiseCast with llvm::bit_cast
Lei Zhang [Fri, 19 Jul 2019 16:51:08 +0000 (09:51 -0700)]
Replace bitwiseCast with llvm::bit_cast

PiperOrigin-RevId: 258986485

5 years agoSuppress compiler warnings regarding unused variables
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

5 years agoWrap op (de)serialization methods in anonymous namespace
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

5 years agoSwitch C++14 std::equal usage to for-loop.
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

5 years agoMove loop dialect tests into separate files - NFC
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

5 years agoAdd missing MLIRDialect dependency for MLIRDialect
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

5 years agoMake SPIR-V spv.EntryPoint and spv.ExecutionMode consistent with SPIR-V spec
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

5 years agoGeneralize implicit terminator into an OpTrait
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

5 years agoUniformize test name - NFC
Nicolas Vasilache [Fri, 19 Jul 2019 13:31:28 +0000 (06:31 -0700)]
Uniformize test name - NFC

PiperOrigin-RevId: 258956693

5 years agoRefactor stripmineSink for AffineForOp - NFC
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

5 years agoUtility function to map a loop on a parametric grid of virtual processors
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

5 years agoUniformize the API for the mlir::tile functions on AffineForOp and loop::ForOp
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

5 years agoAdd support for providing a legality callback for dynamic legality in DialectConversion.
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

5 years ago[spirv] group methods better and improve comments
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

5 years agoPlace generated StandardOps to SPIR-V patterns in anonymous namespace
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

5 years agoNFC: Expose a ConversionPatternRewriter for use with ConversionPatterns.
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

5 years agoAdd an "is_signed" attribute to the quant_ConstFakeQuant op
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

5 years agoMinor cleanup to LangRef, MLIR stands for "Multi-Level IR"
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

5 years agoFix script relative path after moving SPIR-V dialect
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

5 years agoPrint boolean values in ElementsAttr as "true"/"false"
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

5 years agoAdd UnitAttr in OpBase.td.
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

5 years agoAutomatically generate (de)serialization methods for SPIR-V ops
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

5 years agoAdd helper to get flattened tuple types
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

5 years agoSimplify broadcastable traits
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

5 years agoAdd support for parsing/printing the trailing type of a dialect attribute.
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

5 years agoUpdate Contributing.md doc to refer to the developer guide
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

5 years agoRelax Broadcastable trait to only reject instances that are statically incompatible
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

5 years agoRefactor the conversion of block argument types in DialectConversion.
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

5 years agoAdd tests for broadcastable trait
Smit Hinsu [Wed, 17 Jul 2019 21:05:19 +0000 (14:05 -0700)]
Add tests for broadcastable trait

PiperOrigin-RevId: 258637509

5 years agoAdd an initial TestingGuide document to describe testing in MLIR.
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

5 years agoAdd support for explicitly marking dialects and operations as illegal.
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

5 years agoVerify that ReturnOp only appears within the region of a FuncOp.
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

5 years agoMove affine.for and affine.if to ODS
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

5 years agoRefactor DialectConversion to support different conversion modes.
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

5 years agoAdd a TypeIsPred.
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

5 years agoSupport signed and unsigned quantization types
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

5 years agoFix build by making LoopOps depend on StandardOps
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

5 years agoMove shared cpu runner library to Support/JitRunner.
Stephan Herhut [Tue, 16 Jul 2019 12:28:56 +0000 (05:28 -0700)]
Move shared cpu runner library to Support/JitRunner.

PiperOrigin-RevId: 258347825

5 years agoNFC: Move SPIR-V dialect to Dialect/ subdirectory
Lei Zhang [Tue, 16 Jul 2019 12:06:57 +0000 (05:06 -0700)]
NFC: Move SPIR-V dialect to Dialect/ subdirectory

PiperOrigin-RevId: 258345603

5 years agoBetter support for attribute wrapper classes when getting def name
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

5 years agoReplace linalg.for by loop.for
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

5 years agoForward-declare LogicalResult as struct rather than class
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

5 years agoRemove lowerAffineConstructs and lowerControlFlow in favor of providing patterns.
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

5 years agoUpdate 'applyPatternsGreedily' to work on the regions of any operations.
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