[mlir][CF] Split `cf-to-llvm` from `func-to-llvm` (#120580)

Do not run `cf-to-llvm` as part of `func-to-llvm`. This commit fixes
https://github.com/llvm/llvm-project/issues/70982.

This commit ch

[mlir][CF] Split `cf-to-llvm` from `func-to-llvm` (#120580)

Do not run `cf-to-llvm` as part of `func-to-llvm`. This commit fixes
https://github.com/llvm/llvm-project/issues/70982.

This commit changes the way how `func.func` ops are lowered to LLVM.
Previously, the signature of the entire region (i.e., entry block and
all other blocks in the `func.func` op) was converted as part of the
`func.func` lowering pattern.

Now, only the entry block is converted. The remaining block signatures
are converted together with `cf.br` and `cf.cond_br` as part of
`cf-to-llvm`. All unstructured control flow is not converted as part of
a single pass (`cf-to-llvm`). `func-to-llvm` no longer deals with
unstructured control flow.

Also add more test cases for control flow dialect ops.

Note: This PR is in preparation of #120431, which adds an additional
GPU-specific lowering for `cf.assert`. This was a problem because
`cf.assert` used to be converted as part of `func-to-llvm`.

Note for LLVM integration: If you see failures, add
`-convert-cf-to-llvm` to your pass pipeline.

show more ...

[mlir] Fix integration tests after #120548 (#120706)

This should have been part of #120548.

[mlir][bufferization] Remove `finalizing-bufferize` pass (#114154)

The dialect conversion-based bufferization passes have been migrated to
One-Shot Bufferize about two years ago. To clean up the co

[mlir][bufferization] Remove `finalizing-bufferize` pass (#114154)

The dialect conversion-based bufferization passes have been migrated to
One-Shot Bufferize about two years ago. To clean up the code base, this
commit removes the `finalizing-bufferize` pass, one of the few remaining
parts of the old infrastructure. Most bufferization passes have already
been removed.

Note for LLVM integration: If you depend on this pass, migrate to
One-Shot Bufferize or copy the pass to your codebase.

Depends on #114152.

show more ...

[mlir][sparse] add parallelization options to mini pipeline (#104233)

[mlir][vector] Rename LowerVectorToLLVM to ConvertVectorToLLVM (NFC) (#104785)

There was some inconsistency with ConvertVectorToLLVM Pass builder,
files and option names.
This patch aims to move a

[mlir][vector] Rename LowerVectorToLLVM to ConvertVectorToLLVM (NFC) (#104785)

There was some inconsistency with ConvertVectorToLLVM Pass builder,
files and option names.
This patch aims to move all occurences to ConvertVectorToLLVM.

show more ...

[mlir][sparse] expose emit strategy option to mini pipeline (#96238)

[mlir][sparse] add linalg elt-wise fusion to sparsifier pipeline (#90924)

yields better IR in general, and all end-to-end tests pass!

[mlir][linalg] NFC: Use tablegen macro for pass constructors (#82892)

This uses the tablegen macros for generating pass constructors, exposing
pass options for fold-unit-extent-dims and linalg-dete

[mlir][linalg] NFC: Use tablegen macro for pass constructors (#82892)

This uses the tablegen macros for generating pass constructors, exposing
pass options for fold-unit-extent-dims and linalg-detensorize.

Additionally aligns some of the pass namings to their text counterpart.
This includes an API change:

createLinalgGeneralizationPass -> createLinalgGeneralizeNamedOpsPass

show more ...

[mlir][sparse] minor comment edits in sparsifier pipeline (#77000)

Changed all code and comments that used the phrase "sparse compiler" to instead use "sparsifier" (#71875)

The changes in this p.r. mostly center around the tests that use the
flag sparse_compiler (

Changed all code and comments that used the phrase "sparse compiler" to instead use "sparsifier" (#71875)

The changes in this p.r. mostly center around the tests that use the
flag sparse_compiler (also: sparse-compiler).

show more ...

[mlir][sparse][gpu] re-enable all GPU libgen tests (#72185)

Previous change no longer properly used the GPU libgen pass (even though
most tests still passed falling back to CPU). This revision puts

[mlir][sparse][gpu] re-enable all GPU libgen tests (#72185)

Previous change no longer properly used the GPU libgen pass (even though
most tests still passed falling back to CPU). This revision puts the
proper pass order into place. Also bit of a cleanup of CPU codegen vs.
libgen setup.

show more ...

[mlir][sparse] implement sparse_tensor.reorder_coo (#68916)

As a side effect of the change, it also unifies the convertOp
implementation between lib/codegen path.

[mlir][sparse] Extract `StorageSpecifierToLLVMPass` from bufferization pipeline (#68635)

`StorageSpecifierToLLVMPass` does not have to be part of the
bufferization mini pipeline. It can run after t

[mlir][sparse] Extract `StorageSpecifierToLLVMPass` from bufferization pipeline (#68635)

`StorageSpecifierToLLVMPass` does not have to be part of the
bufferization mini pipeline. It can run after the bufferization
pipeline. This is desirable because it keeps the bufferization pipeline
smaller.

Also fix incorrect bufferization API usage: `bufferizeOp` instead of
`bufferizeModuleOp` was used, even though function boundaries were
bufferized.

show more ...

[mlir][gpu][NVPTX] Enable NVIDIA GPU JIT compilation path (#66220)

This patch adds an NVPTX compilation path that enables JIT compilation
on NVIDIA targets. The following modifications were perform

[mlir][gpu][NVPTX] Enable NVIDIA GPU JIT compilation path (#66220)

This patch adds an NVPTX compilation path that enables JIT compilation
on NVIDIA targets. The following modifications were performed:
1. Adding a format field to the GPU object attribute, allowing the
translation attribute to use the correct runtime function to load the
module. Likewise, a dictionary attribute was added to add any possible
extra options.

2. Adding the `createObject` method to `GPUTargetAttrInterface`; this
method returns a GPU object from a binary string.

3. Adding the function `mgpuModuleLoadJIT`, which is only available for
NVIDIA GPUs, as there is no equivalent for AMD.

4. Adding the CMake flag `MLIR_GPU_COMPILATION_TEST_FORMAT` to specify
the format to use during testing.

show more ...

[mlir][gpu] Deprecate gpu::Serialization* passes. (#65857)

Deprecate the `gpu-to-cubin` & `gpu-to-hsaco` passes in favor of the
`TargetAttr` workflow. This patch removes remaining upstream uses of

[mlir][gpu] Deprecate gpu::Serialization* passes. (#65857)

Deprecate the `gpu-to-cubin` & `gpu-to-hsaco` passes in favor of the
`TargetAttr` workflow. This patch removes remaining upstream uses of the
aforementioned passes, including the option to use them in `mlir-opt`. A
future patch will remove these passes entirely.

The passes can be re-enabled in `mlir-opt` by adding the CMake flag: `-DMLIR_ENABLE_DEPRECATED_GPU_SERIALIZATION=1`.

show more ...

Reland [mlir][test][gpu] Migrate CUDA tests to the TargetAttr compilation workflow (llvm#65768)

The revert happened due to a build bot failure that threw 'CUDA_ERROR_UNSUPPORTED_PTX_VERSION'.
The fa

Reland [mlir][test][gpu] Migrate CUDA tests to the TargetAttr compilation workflow (llvm#65768)

The revert happened due to a build bot failure that threw 'CUDA_ERROR_UNSUPPORTED_PTX_VERSION'.
The failure's root cause was a pass using "+ptx76" for compilation and an old CUDA driver
on the bot. This commit relands the patch with "+ptx60".

Original Gh PR: #65768
Original commit message:
Migrate tests referencing `gpu-to-cubin` to the new compilation workflow
using `TargetAttrs`. The `test-lower-to-nvvm` pass pipeline was modified
to use the new compilation workflow to simplify the introduction of
future tests.

The `createLowerGpuOpsToNVVMOpsPass` function was removed, as it didn't
allow for passing all options available in the `ConvertGpuOpsToNVVMOp`
pass.

show more ...

Revert "[mlir][test][gpu] Migrate CUDA tests to the TargetAttr compilation workflow (#65768) (#65848)

This reverts commit d21b67293be15f8a89378e4785d70cc037866406.

[mlir][test][gpu] Migrate CUDA tests to the TargetAttr compilation workflow (#65768)

Migrate tests referencing `gpu-to-cubin` to the new compilation workflow
using `TargetAttrs`. The `test-lower-to

[mlir][test][gpu] Migrate CUDA tests to the TargetAttr compilation workflow (#65768)

Migrate tests referencing `gpu-to-cubin` to the new compilation workflow
using `TargetAttrs`. The `test-lower-to-nvvm` pass pipeline was modified
to use the new compilation workflow to simplify the introduction of
future tests.

The `createLowerGpuOpsToNVVMOpsPass` function was removed, as it didn't
allow for passing all options available in the `ConvertGpuOpsToNVVMOp`
pass.

show more ...

[mlir][memref] Add pass to expand realloc operations, simplify lowering to LLVM

There are two motivations for this change:
1. It considerably simplifies adding support for the realloc operation to t

[mlir][memref] Add pass to expand realloc operations, simplify lowering to LLVM

There are two motivations for this change:
1. It considerably simplifies adding support for the realloc operation to the
new buffer deallocation pass by lowering the realloc such that no
deallocation operation is inserted and the deallocation pass itself can
insert that dealloc
2. The lowering is expressed on a higher level and thus easier to understand,
and the lowerings of the memref operations it is composed of don't have to
be duplicated in the MemRefToLLVM lowering (also see discussion in
https://reviews.llvm.org/D133424)

Reviewed By: springerm

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

show more ...

[mlir][sparse] refactor sparsification and bufferization pass into proper TD pass

Registering the SparsificationAndBufferization into a proper TD pass
has the advantage that it can be invoked and te

[mlir][sparse] refactor sparsification and bufferization pass into proper TD pass

Registering the SparsificationAndBufferization into a proper TD pass
has the advantage that it can be invoked and tested in isolation. This
change also moves some bufferization specific set up from the pipeline
file into the pass file, keeping the logic more locally.

Reviewed By: Peiming

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

show more ...

Revert "[mlir][gpu] Fallback to JIT compilation" "[mlir][gpu] Increase default SM version from 35 to 50" and "[mlir][gpu] Improving Cubin Serialization with ptxas Compiler"

This reverts commit 2e0e

Revert "[mlir][gpu] Fallback to JIT compilation" "[mlir][gpu] Increase default SM version from 35 to 50" and "[mlir][gpu] Improving Cubin Serialization with ptxas Compiler"

This reverts commit 2e0e00ed841951e358a85a871647be9b3a622f51
and reverts commit a6eb40692c795a9cc29266779ceca2e304141114
and reverts commit 585cbe3f639783bf0307b47504acbd205f135310.

15 tests are broken on the mlir-nvidia buildbot:

'cuModuleLoadData(&module, data)' failed with 'CUDA_ERROR_INVALID_SOURCE'
'cuModuleGetFunction(&function, module, name)' failed with 'CUDA_ERROR_INVALID_HANDLE'
'cuLaunchKernel(function, gridX, gridY, gridZ, blockX, blockY, blockZ, smem, stream, params, extra)' failed with 'CUDA_ERROR_INVALID_HANDLE'
'cuModuleUnload(module)' failed with 'CUDA_ERROR_INVALID_HANDLE'

show more ...

[mlir][gpu] Improving Cubin Serialization with ptxas Compiler

This work improves how we compile the generated PTX code using the `ptxas` compiler. Currently, we rely on the driver's jit API to compi

[mlir][gpu] Improving Cubin Serialization with ptxas Compiler

This work improves how we compile the generated PTX code using the `ptxas` compiler. Currently, we rely on the driver's jit API to compile the PTX code. However, this approach has some limitations. It doesn't always produce the same binary output as the ptxas compiler, leading to potential inconsistencies in the generated Cubin files.

This work introduces a significant improvement by directly utilizing the ptxas compiler for PTX compilation. By doing so, we can achieve more consistent and reliable results in generating cubin files. Key Benefits:
- Using the Ptxas compiler directly ensures that the cubin files generated during the build process remain consistent with CUDA compilation using `nvcc` or `clang`.
- Another advantage of this work is that it allows developers to experiment with different ptxas compilers without the need to change the compiler. Performance among ptxas compiler versions are vary, therefore, one can easily try different ptxas compilers.

Reviewed By: nicolasvasilache

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

show more ...

[mlir] Move casting calls from methods to function calls

The MLIR classes Type/Attribute/Operation/Op/Value support
cast/dyn_cast/isa/dyn_cast_or_null functionality through llvm's doCast
functionali

[mlir] Move casting calls from methods to function calls

The MLIR classes Type/Attribute/Operation/Op/Value support
cast/dyn_cast/isa/dyn_cast_or_null functionality through llvm's doCast
functionality in addition to defining methods with the same name.
This change begins the migration of uses of the method to the
corresponding function call as has been decided as more consistent.

Note that there still exist classes that only define methods directly,
such as AffineExpr, and this does not include work currently to support
a functional cast/isa call.

Caveats include:
- This clang-tidy script probably has more problems.
- This only touches C++ code, so nothing that is being generated.

Context:
- https://mlir.llvm.org/deprecation/ at "Use the free function variants
for dyn_cast/cast/isa/…"
- Original discussion at https://discourse.llvm.org/t/preferred-casting-style-going-forward/68443

Implementation:
This first patch was created with the following steps. The intention is
to only do automated changes at first, so I waste less time if it's
reverted, and so the first mass change is more clear as an example to
other teams that will need to follow similar steps.

Steps are described per line, as comments are removed by git:
0. Retrieve the change from the following to build clang-tidy with an
additional check:
https://github.com/llvm/llvm-project/compare/main...tpopp:llvm-project:tidy-cast-check
1. Build clang-tidy
2. Run clang-tidy over your entire codebase while disabling all checks
and enabling the one relevant one. Run on all header files also.
3. Delete .inc files that were also modified, so the next build rebuilds
them to a pure state.
4. Some changes have been deleted for the following reasons:
- Some files had a variable also named cast
- Some files had not included a header file that defines the cast
functions
- Some files are definitions of the classes that have the casting
methods, so the code still refers to the method instead of the
function without adding a prefix or removing the method declaration
at the same time.

```
ninja -C $BUILD_DIR clang-tidy

run-clang-tidy -clang-tidy-binary=$BUILD_DIR/bin/clang-tidy -checks='-*,misc-cast-functions'\
-header-filter=mlir/ mlir/* -fix

rm -rf $BUILD_DIR/tools/mlir/**/*.inc

git restore mlir/lib/IR mlir/lib/Dialect/DLTI/DLTI.cpp\
mlir/lib/Dialect/Complex/IR/ComplexDialect.cpp\
mlir/lib/**/IR/\
mlir/lib/Dialect/SparseTensor/Transforms/SparseVectorization.cpp\
mlir/lib/Dialect/Vector/Transforms/LowerVectorMultiReduction.cpp\
mlir/test/lib/Dialect/Test/TestTypes.cpp\
mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp\
mlir/test/lib/Dialect/Test/TestAttributes.cpp\
mlir/unittests/TableGen/EnumsGenTest.cpp\
mlir/test/python/lib/PythonTestCAPI.cpp\
mlir/include/mlir/IR/
```

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

show more ...

[mlir][bufferization] Make function boundary type convertion logic dynamic.

Having to choose from only static or dynamic layout for all function is limiting.

Differential Revision: https://reviews.

[mlir][bufferization] Make function boundary type convertion logic dynamic.

Having to choose from only static or dynamic layout for all function is limiting.

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

show more ...

[mlir][sparse][gpu] sparse GPU code generator pipeline setup

Reviewed By: Peiming

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