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
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
The function populateStdOpsToSPIRVPatterns appends the conversion
patterns automatically generated from StdOpsToSPIRVConversion.td to a
list of patterns
PiperOrigin-RevId: 259677890
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
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
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
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
- 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>
Closestensorflow/mlir#40
PiperOrigin-RevId: 259442536
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
* 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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
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
Alex Zinenko