[all-commits] [llvm/llvm-project] 0117d8: [MLIR][NVVM] Add optional llc verification for NVV...

[Kirill Vedernikov via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 0117d8ca5e8c37e3a14647875095cc422ee9a686
      https://github.com/llvm/llvm-project/commit/0117d8ca5e8c37e3a14647875095cc422ee9a686
  Author: Kirill Vedernikov <kvedernikov at nvidia.com>
  Date:   2026-09-08 (Tue, 08 Sep 2026)

  Changed paths:
    M mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-collector-b.mlir
    M mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-decompress-b.mlir
    M mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-ti16.mlir
    M mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared.mlir
    M mlir/test/lit.cfg.py

  Log Message:
  -----------
  [MLIR][NVVM] Add optional llc verification for NVVM dialect tests (#219980)

The motivation is to verify that MLIR NVVM dialect operations lower to
LLVM IR that is accepted by the NVPTX backend. Running llc gives these
tests an extra end-to-end check beyond textual IR matching and helps
catch mismatches between MLIR lowering, NVVM intrinsics, and backend
expectations.

A valid LLVM intrinsic call isn't necessarily selectable for every
combination of its immargs and target features. There is a concrete
NVPTX example in
`llvm/test/CodeGen/NVPTX/tcgen05-mma-collector-b-i8-invalid.ll`. The
intrinsic has a valid signature and individually valid immediate values,
but the combination `kind=i8` and `collector_b=lastuse` is not
supported. LLVM IR verification accepts it, while `llc` rejects it with
Cannot select.

The corresponding MLIR verifier does not duplicate every NVPTX
instruction-selection predicate. It currently checks MLIR-level
properties such as the `disable_output_lane size` and the interaction
between `a_shift` and `collector_a` (the examples are from
`tcgen05.mma`), while the backend owns the complete target-dependent
selection rules. We should improve the MLIR verifier whenever an invalid
MLIR operation can be rejected early, but a targeted integration check
also detects disagreement or drift between those two independently
maintained layers.

Another case is target-version drift. A translation test can be extended
with a new intrinsic variant requiring a newer SM or PTX ISA while
continuing to pass its textual `FileCheck`. Since the existing command
has no target contract, the test does not tell us whether all the
generated intrinsics are lowerable for one coherent target. An `llc`
invocation with explicit `-mcpu` and `-mattr` makes that contract
visible. If a test intended for sm_100a/PTX 8.6 starts using
functionality available only for sm_107f/PTX 9.4, the integration check
reports it immediately.

```
// RUN: %if has_llc_binary && host-supports-nvptx %{ mlir-translate --mlir-to-llvmir %s | %llc-verify -march=nvptx64 -mcpu=sm_<version> -mattr=+ptx<version> %}
```
`%llc-verify` finds `llc` as `lit.util.which("llc",
config.llvm_tools_dir)`.



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications