This patch adds support for Calling Convention attribute in LLVM
dialect, including enums, custom syntax and import from LLVM IR.
Additionally fix import of dso_local attribute.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D126161
This is a followup to D126105 to move functions in SparseTensorUtils.cpp to match their locations in SparseTensorUtils.h
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D126106
This change makes the public API of SparseTensorUtils.cpp explicit, whereas before the publicity of these functions was only implicit. Implicit publicity is sufficient for mlir-opt to generate calls to these functions, but it's not enough to enable C/C++ code to call them directly in the usual way (i.e., without going through codegen). Thus, leaving the publicity implicit prevents development of other tools (e.g., microbenchmarks).
In addition this change also marks the functions MLIR_CRUNNERUTILS_EXPORT, which is required by the JIT under certain configurations (albeit not for anything in our test suite).
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D126105
The Transform dialect uses the side effect modeling mechanism to record the
effects of the transform ops on the mapping between Transform IR values and
Payload IR ops. Introduce a checker pass that warns if a Transform IR value is
used after it has been freed (consumed). This pass is mostly intended as a
debugging aid in addition to the verification/assertion mechanisms in the
transform interpreter. It reports all potential use-after-free situations.
The implementation makes a series of simplifying assumptions to be simple and
conservative. A more advanced implementation would rely on the data flow-like
analysis associated with a side-effect resource rather than a value, which is
currently not supported by the analysis infrastructure.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D126381
If we don't specify the result index while matching operand with the
result of certain operation, it's supposed to match all the results of
the operation with the operand. For registered op, it's easy to do that
by either indexing with number or name. For unregistered op, this commit
enables the numeric result indexing for this use case.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D126330
Using 64-bit integer/float type in interface storage classes would
require Int64/Float64 capability, per the Vulkan spec:
```
shaderInt64 specifies whether 64-bit integers (signed and unsigned) are
supported in shader code. If this feature is not enabled, 64-bit integer
types must not be used in shader code. This also specifies whether
shader modules can declare the Int64 capability. Declaring and using
64-bit integers is enabled for all storage classes that SPIR-V allows
with the Int64 capability.
```
This is different from, say, 16-bit element types, where:
```
shaderInt16 specifies whether 16-bit integers (signed and unsigned) are
supported in shader code. If this feature is not enabled, 16-bit integer
types must not be used in shader code. This also specifies whether
shader modules can declare the Int16 capability. However, this only
enables a subset of the storage classes that SPIR-V allows for the Int16
SPIR-V capability: Declaring and using 16-bit integers in the Private,
Workgroup (for non-Block variables), and Function storage classes is
enabled, while declaring them in the interface storage classes (e.g.,
UniformConstant, Uniform, StorageBuffer, Input, Output, and
PushConstant) is not enabled.
```
Reviewed By: hanchung
Differential Revision: https://reviews.llvm.org/D126256
These attributes can carry useful information, e.g., pipelines
might use them to organize and chain patterns.
Reviewed By: hanchung
Differential Revision: https://reviews.llvm.org/D126320
This patch adds support for obtaining a set corresponding to the domain/range
of the relation.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D126326
This patch makes sure that the address dereferences to value in
omp.atomic.write operation.
Reviewed By: kiranchandramohan, peixin
Differential Revision: https://reviews.llvm.org/D126272
Add lowering for cases where the reduction dimension is fully unrolled.
It is common to unroll the reduction dimension, therefore we would want
to lower the contractions to an elementwise vector op in this case.
Differential Revision: https://reviews.llvm.org/D126120
TOSAs depthwise_conv2d operation includes a reshape to include the implicit x1 dimension.
Reviewed By: rsuderman
Differential Revision: https://reviews.llvm.org/D126212
AsyncCopyOp lowering converted "size in elements" to "size in bytes"
assuming the element type size is at least one byte. This removes
that restriction, allowing for types such as i4 and b1 to be handled
correctly.
Differential Revision: https://reviews.llvm.org/D125838
This changes adds missing support for the i4 data type. Tests are added
to ensure proper lowering of an nvgpu.mma.sync operation targeting the
16x8x64xi4 and 16x8x32xi4 MMA variants in the NVVM dialect.
Differential Revision: https://reviews.llvm.org/D126092
Also fixes integration of the pass into One-Shot Bufferize and adds additional test cases.
BufferResultsToOutParams can be used with "identity-layout-map" and "fully-dynamic-layout-map". "infer-layout-map" is not supported.
Differential Revision: https://reviews.llvm.org/D125636
No longer pass static dim sizes as an attribute. This was redundant and required extra checks in the verifier. This change also makes the op symmetrical to memref::AllocOp.
Differential Revision: https://reviews.llvm.org/D126178
This patch updates asserts in IntegerRelation::isEqual and
IntegerRelation::isCompatible to allow these functions when number of
local identifiers are different. This change is done to reflect the
algorithmic changes done before this patch.
One of the ShuffleVectorOp::build functions checks if the incoming
vector operands is scalable vector by casting its type to
mlir::VectorType first. However, in some cases the operand is not
necessarily mlir::VectorType (e.g. it might be a LLVMVectorType).
This patch fixes this issue by using the dedicated
`LLVM::isScalableVectorType` function to determine if the incoming
vector is scalable vector or not.
Differential Revision: https://reviews.llvm.org/D125818
Add support for translating from llvm::Select, llvm::FNeg, and llvm::Unreachable.
This patch also cleans up (NFC) the opcode map for simple instructions and
adds `// clang-format off/on` comments to prevent those lines from being
churned by clang-format between commits.
Differential Revision: https://reviews.llvm.org/D125817
This change adds a new op `alloc_tensor` to the bufferization dialect. During bufferization, this op is always lowered to a buffer allocation (unless it is "eliminated" by a pre-processing pass). It is useful to have such an op in tensor land, because it allows users to model tensor SSA use-def chains (which drive bufferization decisions) and because tensor SSA use-def chains can be analyzed by One-Shot Bufferize, while memref values cannot.
This change also replaces all uses of linalg.init_tensor in bufferization-related code with bufferization.alloc_tensor.
linalg.init_tensor and bufferization.alloc_tensor are similar, but the purpose of the former one is just to carry a shape. It does not indicate a memory allocation.
linalg.init_tensor is not suitable for modelling SSA use-def chains for bufferization purposes, because linalg.init_tensor is marked as not having side effects (in contrast to alloc_tensor). As such, it is legal to move linalg.init_tensor ops around/CSE them/etc. This is not desirable for alloc_tensor; it represents an explicit buffer allocation while still in tensor land and such allocations should not suddenly disappear or get moved around when running the canonicalizer/CSE/etc.
BEGIN_PUBLIC
No public commit message needed for presubmit.
END_PUBLIC
Differential Revision: https://reviews.llvm.org/D126003
This changes adds the option to lower to NvGpu dialect ops during the
VectorToGPU convsersion pass. Because this transformation reuses
existing VectorToGPU logic, a seperate VectorToNvGpu conversion pass is
not created. The option `use-nvgpu` is added to the VectorToGPU pass.
When this is true, the pass will attempt to convert slices rooted at
`vector.contract` operations into `nvgpu.mma.sync` ops, and
`vector.transfer_read` ops are converted to either `nvgpu.ldmatrix` or
one or more `vector.load` operations. The specific data loaded will
depend on the thread id within a subgroup (warp). These index
calculations depend on data type and shape of the MMA op
according to the downstream PTX specification. The code for supporting
these details is separated into `NvGpuSupport.cpp|h`.
Differential Revision: https://reviews.llvm.org/D122940
For the hypothetical "a.b.c" op printed within a region that declares "a" as
the default dialect, MLIR would currently elide the "a." prefix and only print
"b.c". However, this becomes ambiguous while parsing as "b.c" may be exist as
the "c" op in the "b" dialect. If it does not, the parsing currently fails. Do
not elide the default dialect if the op name contains further dots to avoid the
ambiguity.
See https://discourse.llvm.org/t/dropping-dialect-prefix-for-ops-with-multiple-dots-in-the-name/62562
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D125975
By closing over the `rank` itself rather than `this`, we save a method call on each iteration. A minor optimization, but one that adds up.
Depends On D126016
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D126019
Alexander Batashev