summaryrefslogtreecommitdiffstats
path: root/mlir/lib
Commit message (Collapse)AuthorAgeFilesLines
...
* Allow std.constant to hold a boolean value.River Riddle2019-07-221-5/+5
| | | | | | This was an oversight in the original implementation, std.constant already supports IntegerAttr just not BoolAttr. PiperOrigin-RevId: 259467710
* Introduce parser library method to parse list of region argumentsUday Bondhugula2019-07-222-8/+28
| | | | | | | | | | | | | | | | | - 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
* Emit an error for missing '[' when parsing an AffineMapOfSSAIds.River Riddle2019-07-221-1/+1
| | | | | | Fixes tensorflow/mlir#51 PiperOrigin-RevId: 259415034
* (De)serialize composite spv.constantLei Zhang2019-07-222-5/+382
| | | | | | | This CL covers the case of composite spv.constant. We encode/decode them into/from OpConstantComposite/OpConstantNull. PiperOrigin-RevId: 259394700
* NFC: Update the LoopToStd conversion patterns to use RewritePattern instead ↵River Riddle2019-07-221-35/+23
| | | | | | | | of ConversionPattern. These patterns don't require type changes so they don't need to be using ConversionPattern. PiperOrigin-RevId: 259393151
* Update style/clang-format (NFC).Jacques Pienaar2019-07-226-74/+95
| | | | | | Update to be consistent & so that future save + clang-format workflows don't introduce extra changes. PiperOrigin-RevId: 259361174
* (De)serialize float scalar spv.constantLei Zhang2019-07-222-2/+77
| | | | | | This CL adds support for float scalar spv.constant in (de)serialization. PiperOrigin-RevId: 259311776
* (De)serialize bool and integer scalar spv.constantLei Zhang2019-07-223-35/+330
| | | | | | | | | | | | | | | | | | | | | 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
* [spirv] NFC: adjust `encode*` function signatures in SerializerLei Zhang2019-07-221-26/+29
| | | | | | | | | * 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
* [spirv] Remove one level of indirection: processOp to processOpImplLei Zhang2019-07-222-11/+2
| | | | | | | | | | | 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
* Refactor LoopParametricTiling as a test pass - NFCNicolas Vasilache2019-07-223-153/+76
| | | | | | This CL moves LoopParametricTiling into test/lib as a pass for purely testing purposes. PiperOrigin-RevId: 259300264
* Refactor region type signature conversion to be explicit via patterns.River Riddle2019-07-203-210/+221
| | | | | | 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
* Add (de)serialization of EntryPointOp and ExecutionModeOpMahesh Ravishankar2019-07-202-38/+261
| | | | | | | | | 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
* Ensure that DenseElementAttr data is 64-bit aligned.River Riddle2019-07-201-7/+14
| | | | | | 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
* Merge TypeUtilities library into the IR libraryAlex Zinenko2019-07-202-11/+2
| | | | | | | | | | 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
* [spirv] Avoid printing duplicate trailing typeLei Zhang2019-07-191-2/+4
| | | | | | | | | 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
* Replace bitwiseCast with llvm::bit_castLei Zhang2019-07-191-12/+5
| | | | PiperOrigin-RevId: 258986485
* Wrap op (de)serialization methods in anonymous namespaceLei Zhang2019-07-192-0/+4
| | | | | | | | | | | | 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
* Switch C++14 std::equal usage to for-loop.Jacques Pienaar2019-07-191-2/+6
| | | | | | 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
* Make SPIR-V spv.EntryPoint and spv.ExecutionMode consistent with SPIR-V specMahesh Ravishankar2019-07-193-28/+6
| | | | | | | | | | | | | 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
* Generalize implicit terminator into an OpTraitAlex Zinenko2019-07-192-76/+12
| | | | | | | | | | | | | 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
* Refactor stripmineSink for AffineForOp - NFCNicolas Vasilache2019-07-191-34/+9
| | | | | | More moving less cloning. PiperOrigin-RevId: 258947575
* Utility function to map a loop on a parametric grid of virtual processorsNicolas Vasilache2019-07-191-0/+22
| | | | | | | | | | | | | | | | | | | | | | 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
* Uniformize the API for the mlir::tile functions on AffineForOp and loop::ForOpNicolas Vasilache2019-07-194-90/+179
| | | | | | | | | | | 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
* Add support for providing a legality callback for dynamic legality in ↵River Riddle2019-07-191-14/+42
| | | | | | | | | | | | | | | | | | 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
* [spirv] group methods better and improve commentsLei Zhang2019-07-192-253/+310
| | | | | | | | 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
* Place generated StandardOps to SPIR-V patterns in anonymous namespaceLei Zhang2019-07-191-4/+2
| | | | | | This avoids polluting the mlir namespace. PiperOrigin-RevId: 258826497
* NFC: Expose a ConversionPatternRewriter for use with ConversionPatterns.River Riddle2019-07-194-262/+360
| | | | | | 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
* Add an "is_signed" attribute to the quant_ConstFakeQuant opFeng Liu2019-07-192-4/+10
| | | | | | | | | | | 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
* Print boolean values in ElementsAttr as "true"/"false"Lei Zhang2019-07-191-1/+4
| | | | | | | 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
* Automatically generate (de)serialization methods for SPIR-V opsMahesh Ravishankar2019-07-195-39/+123
| | | | | | | | | | | | | | | | 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
* Add helper to get flattened tuple typesGeoffrey Martin-Noble2019-07-191-0/+6
| | | | | | 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
* Add support for parsing/printing the trailing type of a dialect attribute.River Riddle2019-07-194-15/+26
| | | | | | | | | | 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
* Relax Broadcastable trait to only reject instances that are statically ↵Smit Hinsu2019-07-191-10/+36
| | | | | | | | | | | | 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
* Refactor the conversion of block argument types in DialectConversion.River Riddle2019-07-195-307/+286
| | | | | | | | 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
* Add support for explicitly marking dialects and operations as illegal.River Riddle2019-07-191-16/+36
| | | | | | | | 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
* Verify that ReturnOp only appears within the region of a FuncOp.River Riddle2019-07-161-3/+3
| | | | | | The invariants of ReturnOp are directly tied to FuncOp, making ReturnOp invalid in any other context. PiperOrigin-RevId: 258421200
* Move affine.for and affine.if to ODSNicolas Vasilache2019-07-167-104/+91
| | | | | | As the move to ODS is made, body and region names across affine and loop dialects are uniformized. PiperOrigin-RevId: 258416590
* Refactor DialectConversion to support different conversion modes.River Riddle2019-07-165-66/+174
| | | | | | | | | | | 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
* Support signed and unsigned quantization typesFeng Liu2019-07-161-8/+13
| | | | | | | | 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
* Fix build by making LoopOps depend on StandardOpsAlex Zinenko2019-07-161-1/+1
| | | | | | LoopOps needs the definition ConstantIndexOp in the verifier of loop::ForOp. PiperOrigin-RevId: 258355329
* Move shared cpu runner library to Support/JitRunner.Stephan Herhut2019-07-162-0/+345
| | | | PiperOrigin-RevId: 258347825
* NFC: Move SPIR-V dialect to Dialect/ subdirectoryLei Zhang2019-07-1615-30/+30
| | | | PiperOrigin-RevId: 258345603
* Better support for attribute wrapper classes when getting def nameLei Zhang2019-07-161-2/+11
| | | | | | | | | | | | | | | | | | | | | 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
* Replace linalg.for by loop.forNicolas Vasilache2019-07-168-219/+63
| | | | | | | 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
* Remove lowerAffineConstructs and lowerControlFlow in favor of providing ↵River Riddle2019-07-164-27/+28
| | | | | | | | 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
* Update 'applyPatternsGreedily' to work on the regions of any operations.River Riddle2019-07-161-20/+24
| | | | | | 'applyPatternsGreedily' is a useful utility outside of just function regions. PiperOrigin-RevId: 258182937
* Refactor the traversal of operations to Convert in DialectConversion.River Riddle2019-07-161-99/+79
| | | | | | | | 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
* Fix opt build (unused variable in Linalg).Andy Davis2019-07-161-0/+1
| | | | PiperOrigin-RevId: 258168108
* Linalg Utils: use Doxygen comments where appropriateAlex Zinenko2019-07-161-3/+3
| | | | PiperOrigin-RevId: 258160982
OpenPOWER on IntegriCloud