summaryrefslogtreecommitdiffstats
path: root/mlir/test
Commit message (Collapse)AuthorAgeFilesLines
...
* Fix linalg_matmul_impl interfacing with sgemmNicolas Vasilache2019-07-262-2/+2
| | | | | | 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
* Use "standard" load and stores in LowerVectorTransfersNicolas Vasilache2019-07-261-4/+4
| | | | | | | 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
* Automated rollback of commit fc194a4f22fe53f46729821d9c4a993fe200facfMehdi Amini2019-07-252-97/+0
| | | | PiperOrigin-RevId: 260037115
* [spirv] Add AccessChainOp operation.Denis Khalikov2019-07-251-0/+119
| | | | | | | | | AccessChainOp creates a pointer into a composite object that can be used with OpLoad and OpStore. Closes tensorflow/mlir#52 PiperOrigin-RevId: 260035676
* Add support for hexadecimal float literalsAlex Zinenko2019-07-252-0/+97
| | | | | | | | | | | | | 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
* Add support for an analysis mode to DialectConversion.River Riddle2019-07-252-5/+59
| | | | | | | | 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
* Fix backward slice computation to iterate through known control flowNicolas Vasilache2019-07-251-1/+54
| | | | | | | | | 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
* Move GPU dialect to {lib,include/mlir}/DialectAlex Zinenko2019-07-254-0/+0
| | | | | | | 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
* Move SPIRV dialect tests under test/DialectAlex Zinenko2019-07-2411-0/+0
| | | | | | | This was overlooked when the dialect code was moved under {lib,include/mlir}/Dialect. NFC. PiperOrigin-RevId: 259801927
* Cleanup slicing test.Nicolas Vasilache2019-07-242-194/+207
| | | | | | Remove hardcoded SSA names and make use of CHECK-LABEL directives. PiperOrigin-RevId: 259767803
* Enable multi-level Linalg fusionNicolas Vasilache2019-07-241-0/+14
| | | | | | | | 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
* Add sitofp to the standard dialectMLIR Team2019-07-233-3/+45
| | | | | | 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
* Affine loop parallelism detection: conservatively handle unknown opsAlex Zinenko2019-07-232-4/+26
| | | | | | | | | | | | | | | | | 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
* ODS: support UnitAttr in Operation definitionsAlex Zinenko2019-07-233-0/+37
| | | | | | | | | | | | | | | 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
* Introduce LLVMFuncOpAlex Zinenko2019-07-231-0/+87
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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
* Allow std.constant to hold a boolean value.River Riddle2019-07-222-3/+9
| | | | | | 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-223-0/+49
| | | | | | | | | | | | | | | | | - 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-0/+9
| | | | | | Fixes tensorflow/mlir#51 PiperOrigin-RevId: 259415034
* (De)serialize composite spv.constantLei Zhang2019-07-221-0/+28
| | | | | | | This CL covers the case of composite spv.constant. We encode/decode them into/from OpConstantComposite/OpConstantNull. PiperOrigin-RevId: 259394700
* (De)serialize float scalar spv.constantLei Zhang2019-07-221-0/+27
| | | | | | 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-221-0/+37
| | | | | | | | | | | | | | | | | | | | | 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
* Refactor LoopParametricTiling as a test pass - NFCNicolas Vasilache2019-07-223-2/+74
| | | | | | This CL moves LoopParametricTiling into test/lib as a pass for purely testing purposes. PiperOrigin-RevId: 259300264
* SingleBlockImplicitTerminator: report the wrong terminator op foundAlex Zinenko2019-07-221-1/+1
| | | | | | | | | | 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
* Refactor region type signature conversion to be explicit via patterns.River Riddle2019-07-202-7/+10
| | | | | | 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-204-1/+46
| | | | | | | | | 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
* Merge TypeUtilities library into the IR libraryAlex Zinenko2019-07-202-3/+3
| | | | | | | | | | 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/+2
| | | | | | | | | 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
* 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-191-33/+20
| | | | | | | | | | | | | 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-191-3/+6
| | | | | | | | | | | | | 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
* Utility function to map a loop on a parametric grid of virtual processorsNicolas Vasilache2019-07-193-0/+106
| | | | | | | | | | | | | | | | | | | | | | 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-191-6/+73
| | | | | | | | | | | 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/+10
| | | | | | | | | | | | | | | | | | 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
* NFC: Expose a ConversionPatternRewriter for use with ConversionPatterns.River Riddle2019-07-191-8/+12
| | | | | | 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-191-1/+46
| | | | | | | | | | | 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/+1
| | | | | | | 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-191-1/+21
| | | | | | Note that UnitAttr cannot be used for op definition yet. PiperOrigin-RevId: 258693338
* Automatically generate (de)serialization methods for SPIR-V opsMahesh Ravishankar2019-07-192-0/+27
| | | | | | | | | | | | | | | | 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
* 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-191-0/+6
| | | | | | | | | | 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-15/+45
| | | | | | | | | | | | 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-193-6/+56
| | | | | | | | 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 support for explicitly marking dialects and operations as illegal.River Riddle2019-07-193-1/+11
| | | | | | | | 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-0/+10
| | | | | | The invariants of ReturnOp are directly tied to FuncOp, making ReturnOp invalid in any other context. PiperOrigin-RevId: 258421200
* Refactor DialectConversion to support different conversion modes.River Riddle2019-07-161-2/+2
| | | | | | | | | | | 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-162-0/+17
| | | | | | | | 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
* Replace linalg.for by loop.forNicolas Vasilache2019-07-169-160/+102
| | | | | | | 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
OpenPOWER on IntegriCloud