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
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
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
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
We only verify broadcastable trait verifier and don't care about mutations so removed all CHECK statements and FileCheck invocation.
PiperOrigin-RevId: 258662882
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
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
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
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
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
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
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
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
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
Lei Zhang