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

[llvmlistbot at llvm.org]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir

@llvm/pr-subscribers-mlir-llvm

Author: Kirill Vedernikov (kvederni)

<details>
<summary>Changes</summary>

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)`

---
Full diff: https://github.com/llvm/llvm-project/pull/219980.diff


5 Files Affected:

- (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-collector-b.mlir (+1) 
- (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-decompress-b.mlir (+1) 
- (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared-ti16.mlir (+1) 
- (modified) mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-shared.mlir (+1) 
- (modified) mlir/test/lit.cfg.py (+21) 


``````````diff
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)
 
 

``````````

</details>


https://github.com/llvm/llvm-project/pull/219980