[Dialect] Migrate away from PointerUnion::{is,get} (NFC) (#120818)

Note that PointerUnion::{is,get} have been soft deprecated in
PointerUnion.h:

// FIXME: Replace the uses of is(), get() and d

[Dialect] Migrate away from PointerUnion::{is,get} (NFC) (#120818)

Note that PointerUnion::{is,get} have been soft deprecated in
PointerUnion.h:

// FIXME: Replace the uses of is(), get() and dyn_cast() with
// isa<T>, cast<T> and the llvm::dyn_cast<T>

I'm not touching PointerUnion::dyn_cast for now because it's a bit
complicated; we could blindly migrate it to dyn_cast_if_present, but
we should probably use dyn_cast when the operand is known to be
non-null.

show more ...

[mlir] Construct SmallVector with ArrayRef (NFC) (#101896)

[mlir] split transform interfaces into a separate library (#85221)

Transform interfaces are implemented, direction or via extensions, in
libraries belonging to multiple other dialects. Those dialec

[mlir] split transform interfaces into a separate library (#85221)

Transform interfaces are implemented, direction or via extensions, in
libraries belonging to multiple other dialects. Those dialects don't
need to depend on the non-interface part of the transform dialect, which
includes the growing number of ops and transitive dependency footprint.

Split out the interfaces into a separate library. This in turn requires
flipping the dependency from the interface on the dialect that has crept
in because both co-existed in one library. The interface shouldn't
depend on the transform dialect either.

As a consequence of splitting, the capability of the interpreter to
automatically walk the payload IR to identify payload ops of a certain
kind based on the type used for the entry point symbol argument is
disabled. This is a good move by itself as it simplifies the interpreter
logic. This functionality can be trivially replaced by a
`transform.structured.match` operation.

show more ...

Apply clang-tidy fixes for readability-identifier-naming in Utils.cpp (NFC)

Apply clang-tidy fixes for performance-unnecessary-value-param in Utils.cpp (NFC)

Apply clang-tidy fixes for bugprone-macro-parentheses in Utils.cpp (NFC)

[mlir][transforms] Revamp the implementation of mapping loops to GPUs

This revision significantly simplifies the specification and implementation of mapping loops to GPU ids.

Each type of mapping (

[mlir][transforms] Revamp the implementation of mapping loops to GPUs

This revision significantly simplifies the specification and implementation of mapping loops to GPU ids.

Each type of mapping (block, warpgroup, warp, thread) now comes with 2 mapping modes:
1. a 3-D "grid-like" mode, subject to alignment considerations on threadIdx.x, on which predication
may occur on a per-dimension 3-D sub-rectangle basis.
2. a n-D linearized mode, on which predication may only occur on a linear basis.

In the process, better size and alignment requirement inference are introduced along with improved runtime verification messages.

The `warp_dims` attribute was deemed confusing and is removed from the transform in favor of better size inference.

Differential Revision: https://reviews.llvm.org/D155941

show more ...

[mlir][linalg] NFC - Move some utils in preparation for revamping mapping of scf.forall