[Mlir-commits] [mlir] [MLIR][NVVM] Add optional llc verification for NVVM dialect tests (PR #219980)

[Kirill Vedernikov]

https://github.com/kvederni created https://github.com/llvm/llvm-project/pull/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.
```
// RUN: %if llc && host-supports-nvptx %{ mlir-translate --mlir-to-llvmir %s | %llc-verify -march=nvptx64 -mcpu=sm_<version> -mattr=+ptx<version> %}
```
`%llc-verify` finds `llc` in this order:
1. Lit parameter override: `./build/bin/llvm-lit -D llc=/path/to/llc ...`
2. Environment variable: `LLVM_LLC_EXECUTABLE=/path/to/llc ./build/bin/llvm-lit ...`
3. In-tree build tools directory: `lit.util.which("llc", config.llvm_tools_dir)`

>From b7b2b041aa6f8525766b2989977a2be487ec3c60 Mon Sep 17 00:00:00 2001
From: Kirill Vedernikov <kvedernikov at nvidia.com>
Date: Mon, 31 Aug 2026 15:25:16 +0200
Subject: [PATCH] [MLIR][NVVM] Add optional llc verification for NVVM dialect
 tests

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.
---
 .../nvvm/tcgen05-mma-shared-collector-b.mlir  |  1 +
 .../nvvm/tcgen05-mma-shared-decompress-b.mlir |  1 +
 .../LLVMIR/nvvm/tcgen05-mma-shared-ti16.mlir  |  1 +
 .../LLVMIR/nvvm/tcgen05-mma-shared.mlir       |  1 +
 mlir/test/lit.cfg.py                          | 21 +++++++++++++++++++
 5 files changed, 25 insertions(+)

diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-collector-b.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-collector-b.mlir
index 6c11138c2cefb..fc1b6c0d54442 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-collector-b.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-collector-b.mlir
@@ -1,4 +1,5 @@
 // RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s
+// RUN: %if llc && host-supports-nvptx %{ mlir-translate --mlir-to-llvmir %s | %llc-verify -march=nvptx64 -mcpu=sm_107f -mattr=+ptx94 %}
 
 // CHECK-LABEL: @nvvm_tcgen05_mma_collector_b
 llvm.func @nvvm_tcgen05_mma_collector_b(%d_tmem : !llvm.ptr<6>, %a_desc: i64, %b_desc: i64, %idesc: i32, %enable_input_d: i1) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-decompress-b.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-decompress-b.mlir
index 2a568f5961e47..1697fc397af68 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-decompress-b.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-decompress-b.mlir
@@ -1,4 +1,5 @@
 // RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s
+// RUN: %if llc && host-supports-nvptx %{ mlir-translate --mlir-to-llvmir %s | %llc-verify -march=nvptx64 -mcpu=sm_107f -mattr=+ptx94 %}
 
 // CHECK-LABEL: @nvvm_tcgen05_mma_shared_decompress_b_cta_1
 llvm.func @nvvm_tcgen05_mma_shared_decompress_b_cta_1(
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-ti16.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-ti16.mlir
index 37915177e5ce9..699ff602e0b9d 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-ti16.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-ti16.mlir
@@ -1,4 +1,5 @@
 // RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s
+// RUN: %if llc && host-supports-nvptx %{ mlir-translate --mlir-to-llvmir %s | %llc-verify -march=nvptx64 -mcpu=sm_107f -mattr=+ptx94 %}
 
 // CHECK-LABEL: @nvvm_tcgen05_mma_ti16
 llvm.func @nvvm_tcgen05_mma_ti16(%d_tmem : !llvm.ptr<6>, %a_desc: i64, %b_desc: i64, %idesc: i32, %enable_input_d: i1) {
diff --git a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared.mlir b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared.mlir
index 5c259ae7ce18a..c7cffbf19e98f 100644
--- a/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared.mlir
@@ -1,4 +1,5 @@
 // RUN: mlir-translate --mlir-to-llvmir %s | FileCheck %s
+// RUN: %if llc && host-supports-nvptx %{ mlir-translate --mlir-to-llvmir %s | %llc-verify -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 %}
 
 // CHECK-LABEL: @nvvm_tcgen05_mma_cta_1
 llvm.func @nvvm_tcgen05_mma_cta_1(%d_tmem : !llvm.ptr<6>, %a_desc: i64, %adesc: i64, %b_desc: i64, %idesc: i32, %enable_input_d: i1) {
diff --git a/mlir/test/lit.cfg.py b/mlir/test/lit.cfg.py
index 6801dbac1ced9..6ecbd5de40fb9 100644
--- a/mlir/test/lit.cfg.py
+++ b/mlir/test/lit.cfg.py
@@ -316,6 +316,27 @@ def find_real_python_interpreter():
 else:
     tools.extend(["mlir-opt"])
 
+def enable_llc(llc_executable):
+    config.available_features.add("llc")
+    tools.extend(
+        [
+            ToolSubst("%llc", llc_executable),
+            ToolSubst(
+                "%llc-verify",
+                f"{llc_executable} -verify-machineinstrs -o /dev/null",
+            ),
+        ]
+    )
+
+
+llc_executable = (
+    lit_config.params.get("llc")
+    or os.environ.get("LLVM_LLC_EXECUTABLE", None)
+    or lit.util.which("llc", config.llvm_tools_dir)
+)
+if llc_executable:
+    enable_llc(llc_executable)
+
 llvm_config.add_tool_substitutions(tools, tool_dirs)