summaryrefslogtreecommitdiffstats
path: root/mlir
Commit message (Collapse)AuthorAgeFilesLines
...
* Refactor region type signature conversion to be explicit via patterns.River Riddle2019-07-2010-267/+276
| | | | | | 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-208-45/+315
| | | | | | | | | 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-206-15/+6
| | | | | | | | | | 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
* Fix a comment about ShapedType::getNumElements()MLIR Team2019-07-191-1/+1
| | | | PiperOrigin-RevId: 259057303
* [spirv] Avoid printing duplicate trailing typeLei Zhang2019-07-192-4/+6
| | | | | | | | | 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
* Suppress compiler warnings regarding unused variablesLei Zhang2019-07-191-2/+2
| | | | | | | 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
* 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
* Move loop dialect tests into separate files - NFCAlex Zinenko2019-07-194-160/+166
| | | | | | | This was overlooked when moving out loop operations from Standard to a separate dialect. PiperOrigin-RevId: 258970115
* Add missing MLIRDialect dependency for MLIRDialectLei Zhang2019-07-191-0/+1
| | | | | | | Tests for the Broadcastable trait was added in a previous CL, which makes MLIRDialect depend on MLIRDialect now. PiperOrigin-RevId: 258967332
* Make SPIR-V spv.EntryPoint and spv.ExecutionMode consistent with SPIR-V specMahesh Ravishankar2019-07-196-81/+34
| | | | | | | | | | | | | 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-198-104/+100
| | | | | | | | | | | | | 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
* Uniformize test name - NFCNicolas Vasilache2019-07-191-0/+0
| | | | PiperOrigin-RevId: 258956693
* 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-195-0/+161
| | | | | | | | | | | | | | | | | | | | | | 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-196-101/+266
| | | | | | | | | | | 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-194-36/+116
| | | | | | | | | | | | | | | | | | 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-193-255/+312
| | | | | | | | 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-1912-306/+469
| | | | | | 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-194-6/+59
| | | | | | | | | | | 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
* Minor cleanup to LangRef, MLIR stands for "Multi-Level IR"Mehdi Amini2019-07-191-6/+4
| | | | PiperOrigin-RevId: 258798577
* Fix script relative path after moving SPIR-V dialectLei Zhang2019-07-193-3/+3
| | | | PiperOrigin-RevId: 258786729
* Print boolean values in ElementsAttr as "true"/"false"Lei Zhang2019-07-192-2/+5
| | | | | | | 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
* Add UnitAttr in OpBase.td.Jing Pu2019-07-192-1/+29
| | | | | | Note that UnitAttr cannot be used for op definition yet. PiperOrigin-RevId: 258693338
* Automatically generate (de)serialization methods for SPIR-V opsMahesh Ravishankar2019-07-1913-78/+493
| | | | | | | | | | | | | | | | 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-193-8/+12
| | | | | | 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
* Simplify broadcastable traitsSmit Hinsu2019-07-191-25/+4
| | | | | | We only verify broadcastable trait verifier and don't care about mutations so removed all CHECK statements and FileCheck invocation. PiperOrigin-RevId: 258662882
* Add support for parsing/printing the trailing type of a dialect attribute.River Riddle2019-07-197-22/+43
| | | | | | | | | | 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
* Update Contributing.md doc to refer to the developer guideMehdi Amini2019-07-191-7/+2
| | | | | | | | 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
* Relax Broadcastable trait to only reject instances that are statically ↵Smit Hinsu2019-07-193-29/+85
| | | | | | | | | | | | 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-1916-390/+401
| | | | | | | | 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 tests for broadcastable traitSmit Hinsu2019-07-193-0/+150
| | | | PiperOrigin-RevId: 258637509
* Add an initial TestingGuide document to describe testing in MLIR.River Riddle2019-07-192-0/+175
| | | | | | 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
* Add support for explicitly marking dialects and operations as illegal.River Riddle2019-07-195-18/+73
| | | | | | | | 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-163-4/+18
| | | | | | 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-1610-343/+329
| | | | | | 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-1611-95/+232
| | | | | | | | | | | 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
* Add a TypeIsPred.Jacques Pienaar2019-07-163-0/+24
| | | | | | | | 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
* Support signed and unsigned quantization typesFeng Liu2019-07-162-11/+17
| | | | | | | | 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-167-40/+96
| | | | PiperOrigin-RevId: 258347825
* NFC: Move SPIR-V dialect to Dialect/ subdirectoryLei Zhang2019-07-1626-51/+51
| | | | PiperOrigin-RevId: 258345603
* Better support for attribute wrapper classes when getting def nameLei Zhang2019-07-163-16/+38
| | | | | | | | | | | | | | | | | | | | | 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-1622-493/+183
| | | | | | | 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
* Forward-declare LogicalResult as struct rather than classAlex Zinenko2019-07-161-1/+1
| | | | | | | Windows builds are broken by a class/struct mismatch between a forward-declaration of LogicalResult and its definition. PiperOrigin-RevId: 258320420
* Remove lowerAffineConstructs and lowerControlFlow in favor of providing ↵River Riddle2019-07-168-55/+60
| | | | | | | | 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-163-25/+31
| | | | | | 'applyPatternsGreedily' is a useful utility outside of just function regions. PiperOrigin-RevId: 258182937
OpenPOWER on IntegriCloud