The revision renames some methods of the Importer and changes
the error handling to be closer the ModuleTranslation. In particular,
processValue -> lookupValue and processType -> convertType
now fail if the translation fails (instead of returning an error),
which simplifies the error handling.
The revision prepares a follow up commit that will import
LLVMIR intrinsics using tablegen.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D135349
Inserting a tensor into an equivalent tensor is a no-op after bufferization. No alloc is needed.
Differential Revision: https://reviews.llvm.org/D132662
The ConvertVectorToGpu pass implementation contained a small private
support library for performing various calculations during conversion
between `vector` and `nvgpu.mma.sync` and `nvgpu.ldmatrix` operations.
The support library is moved under `Dialect/NVGPU/Utils` because the
functions have wider utility. Some documentation comments are added or
improved.
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D135303
This differential adjusts the numeric values for DimLevelType values: using the low-order two bits for recording the "No" and "Nu" properties, and the high-order bits for the formats per se. (The choice of encoding may seem a bit peculiar, since the bits are mapped to negative properties rather than positive properties. But this was done in order to preserve the collation order of DimLevelType values. If we don't care about collation order, then we may prefer to flip the semantics of the property bits, so that they're less surprising to readers.)
Using distinguished bits for the properties and formats enables faster implementation for the predicates detecting those properties/formats, which matters because this is in the runtime library itself (rather than on the codegen side of things). This differential pushes through the changes to the enum values, and optimizes the basic predicates. However it does not optimize all the places where we check compound predicates (e.g., "is compressed or singleton"), to help reduce rebasing conflict with D134933. Those optimizations will be done after this differential and D134933 are landed.
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D135004
If the scaling factor is by 1 with no offset or border, then the
resize is a no-op.
Reviewed By: dcaballe
Differential Revision: https://reviews.llvm.org/D135329
Handle more cases of singleton DLT including direct sparse2sparse conversion. (Followup to D134096)
Depends On D134926
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D134933
Attribute stride and shift are removed, and has new scale and border.
Signed-off-by: TatWai Chong <tatwai.chong@arm.com>
Change-Id: I6cdbeb3978f5ee540bc6cf59eb7c273eb0131430
Reviewed By: rsuderman
Differential Revision: https://reviews.llvm.org/D131629
The current splicing behavior dates back to when all blocks had terminators,
so we would "helpfully" splice before the terminator. This doesn't make sense
anymore, and leads to somewhat unexpected results when parsing multiple
pieces of IR into the same block.
Differential Revision: https://reviews.llvm.org/D135096
Motivation: we have lowering pipeline based on upstream gpu and spirv dialects and and we are using host shared gpu memory to transfer data between host and device.
Add `host_shared` flag to `gpu.alloc` to distinguish between shared and device-only gpu memory allocations.
Differential Revision: https://reviews.llvm.org/D133533
This patch moves the 'printOp' functionality to the public API of
AsmPrinter and rename it to 'printCustomOrGenericOp'. No 'parseOp'
is needed at this time as existing APIs are able to parse operations
producing results where results are omitted in the textual form
(the LHS of an operation is redundant when it comes to building the
operation itself as it only contains the result names).
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D135006
UnitAttr is optional but unwrapped builders require it. Make Change onstructing
from bool as required for when not set at moment (for UnitAttr nothing needs to
be constructed, this is true for others here too and can be addressed
together).
Differential Revision: https://reviews.llvm.org/D135058
This patch exposes the method to check if an op can be vectorized or
not for downstream uses. Also adds a check to mark elementwise operations
that have non-vectorizable ops (like `tensor.extract`) as non vectorizable.
Reviewed By: nicolasvasilache, dcaballe, ThomasRaoux
Differential Revision: https://reviews.llvm.org/D135201
This extension to the sparse tensor type system in MLIR
opens up a whole new set of sparse storage schemes, such as
block sparse storage (e.g. BCSR) and ELL (aka jagged diagonals).
This revision merely introduces the type extension and
initial documentation. The actual interpretation of the type
(reading in tensors, lowering to code, etc.) will follow.
Reviewed By: Peiming
Differential Revision: https://reviews.llvm.org/D135206
The test Dialect/Affine/ops.mlir was failing when building with
Visual Studio 2022 version 17.3.5. This was caused by a bad MSVC codegen, when
capturing a `constexpr` in a lambda. The bug was reported to Microsoft, see
differential for more information.
Differential revision: https://reviews.llvm.org/D134227
The GPU transform dialect currently has restrictions and several situations where we can't use transform dialect.
This update includes a method to test a failing cases in GPU transform dialect.
Differential Revision: https://reviews.llvm.org/D135063
(Re-Apply with fixes to clang MicrosoftMangle.cpp)
This is a first step towards high level representation for fp8 types
that have been built in to hardware with near term roadmaps. Like the
BFLOAT16 type, the family of fp8 types are inspired by IEEE-754 binary
floating point formats but, due to the size limits, have been tweaked in
various ways in order to maximally use the range/precision in various
scenarios. The list of variants is small/finite and bounded by real
hardware.
This patch introduces the E5M2 FP8 format as proposed by Nvidia, ARM,
and Intel in the paper: https://arxiv.org/pdf/2209.05433.pdf
As the more conformant of the two implemented datatypes, we are plumbing
it through LLVM's APFloat type and MLIR's type system first as a
template. It will be followed by the range optimized E4M3 FP8 format
described in the paper. Since that format deviates further from the
IEEE-754 norms, it may require more debate and implementation
complexity.
Given that we see two parts of the FP8 implementation space represented
by these cases, we are recommending naming of:
* `F8M<N>` : For FP8 types that can be conceived of as following the
same rules as FP16 but with a smaller number of mantissa/exponent
bits. Including the number of mantissa bits in the type name is enough
to fully specify the type. This naming scheme is used to represent
the E5M2 type described in the paper.
* `F8M<N>F` : For FP8 types such as E4M3 which only support finite
values.
The first of these (this patch) seems fairly non-controversial. The
second is previewed here to illustrate options for extending to the
other known variant (but can be discussed in detail in the patch
which implements it).
Many conversations about these types focus on the Machine-Learning
ecosystem where they are used to represent mixed-datatype computations
at a high level. At that level (which is why we also expose them in
MLIR), it is important to retain the actual type definition so that when
lowering to actual kernels or target specific code, the correct
promotions, casts and rescalings can be done as needed. We expect that
most LLVM backends will only experience these types as opaque `I8`
values that are applicable to some instructions.
MLIR does not make it particularly easy to add new floating point types
(i.e. the FloatType hierarchy is not open). Given the need to fully
model FloatTypes and make them interop with tooling, such types will
always be "heavy-weight" and it is not expected that a highly open type
system will be particularly helpful. There are also a bounded number of
floating point types in use for current and upcoming hardware, and we
can just implement them like this (perhaps looking for some cosmetic
ways to reduce the number of places that need to change). Creating a
more generic mechanism for extending floating point types seems like it
wouldn't be worth it and we should just deal with defining them one by
one on an as-needed basis when real hardware implements a new scheme.
Hopefully, with some additional production use and complete software
stacks, hardware makers will converge on a set of such types that is not
terribly divergent at the level that the compiler cares about.
(I cleaned up some old formatting and sorted some items for this case:
If we converge on landing this in some form, I will NFC commit format
only changes as a separate commit)
Differential Revision: https://reviews.llvm.org/D133823
The actual transformation doesn't support multi-output GenericOps, but
if we encounter one without sparse annotations we can just leave it
alone.
Differential Revision: https://reviews.llvm.org/D135176
If 'order(concurrent)' clause is specified, then the iterations of SIMD loop
can be executed concurrently.
This patch adds support for LLVM IR codegen via OMPIRBuilder for SIMD loop
with 'order(concurrent)' clause. The functionality added to OMPIRBuilder is
similar to the functionality implemented in 'CodeGenFunction::EmitOMPSimdInit'.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D134046
Signed-off-by: Dominik Adamski <dominik.adamski@amd.com>
Simplify affine expressions and maps while exploiting simple range and
step info of any IVs that are operands. This simplification is local,
O(1) and practically useful in several scenarios. Accesses with
floordiv's and mod's where the LHS is non-negative and bounded or is a
known multiple of a constant can often be simplified. This is
implemented as a canonicalization for all affine ops in a generic way:
all affine.load/store, vector_load/store, affine.apply, affine.min/max,
etc. ops.
Eg: For tiled loop nests accessing buffers this way:
affine.for %i = 0 to 1024 step 32 {
affine.for %ii = 0 to 32 {
affine.load [(%i + %ii) floordiv 32, (%i + %ii) mod 32]
}
}
// Note that %i is a multiple of 32 and %ii < 32, hence:
(%i + %ii) floordiv 32 is the same as %i floordiv 32
(%i + %ii) mod 32 is the same as %ii mod 32.
The simplification leads to simpler index/subscript arithmetic for
multi-dimensional arrays and also in turn enables detection of spatial
locality (for vectorization for eg.), temporal locality or loop
invariance for hoisting or scalar replacement.
Differential Revision: https://reviews.llvm.org/D135085
Relax the restriction in the transform dialect interpreter utilities
that expected a payload IR op to be assocaited with at most one
transform IR handle value. This was useful during the initial
bootstrapping to avoid use-after-free error equivalents when a payload
IR op could be erased through one of the handles associated with it and
then accessed through another. It was, however, possible to erase an
ancestor of the payload IR operation in question. The expensive-checks
mode of interpretation is able to detect both cases and has proven
sufficiently robust in debugging use-after-free errors.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D134964
This revision adds GPU transform dialect. It also introduce a prefix such as "transform.gpu" for all ops related to this dialect.
MLIR already had two GPU transform op in linalg. This revision moves these ops into GPUTransformOps. The Ops are as follows:
`transform.structured.map_nested_foreach_thread_to_gpu_blocks` -> `transform.gpu.map_foreach_to_blocks`
This op selects the outermost (toplevel) foreach_thread and parallelize across GPU blocks. It can also generate `gpu_launch`.
`transform.structured.map_nested_foreach_thread_to_gpu_threads` -> `transform.gpu.map_nested_foreach_to_threads`
This op parallelizes nested foreach_thread that are inside `gpu_launch` across GPU threads.
It doesn't add new functionality, but there are some minor refactoring of the code.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D134800
tensor.empty/linalg.init_tensor produces an uninititalized tensor that can be used as a destination operand for destination-style ops (ops that implement `DestinationStyleOpInterface`).
This change makes it possible to implement `TilingInterface` for non-destination-style ops without depending on the Linalg dialect.
RFC: https://discourse.llvm.org/t/rfc-add-tensor-from-shape-operation/65101
Differential Revision: https://reviews.llvm.org/D135129
bixia1