[Mlir-commits] [mlir] 6778f0d - [MLIR][NVVM][Tests]: Update FileCheck primitives (#173252)

[llvmlistbot at llvm.org]

Author: Durgadoss R
Date: 2025-12-23T00:01:25+05:30
New Revision: 6778f0d4834904c036113906e2a82aa6b97ad793

URL: https://github.com/llvm/llvm-project/commit/6778f0d4834904c036113906e2a82aa6b97ad793
DIFF: https://github.com/llvm/llvm-project/commit/6778f0d4834904c036113906e2a82aa6b97ad793.diff

LOG: [MLIR][NVVM][Tests]: Update FileCheck primitives (#173252)

This patch updates a few FileCheck primitives for the TMA test
to use CHECK-PTX-DAG instead of CHECK-PTX to accommodate
a slightly different ordering of BB's.

The dump-ptx integration test fails when the PTX is generated
through nvcc (intermediates) from public toolkit. This patch fixes
it by allowing regex strings from both the backends.

Signed-off-by: Durgadoss R <durgadossr at nvidia.com>

Added: 
    

Modified: 
    mlir/test/Integration/GPU/CUDA/dump-ptx.mlir
    mlir/test/Integration/GPU/CUDA/sm90/transform-dialect/tma_load_64x8_8x128_noswizzle-transform.mlir

Removed: 
    


################################################################################
diff  --git a/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir b/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir
index efffcaaf23b2e..5b40946ce3537 100644
--- a/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir
+++ b/mlir/test/Integration/GPU/CUDA/dump-ptx.mlir
@@ -2,7 +2,7 @@
 // RUN:  | mlir-opt -gpu-lower-to-nvvm-pipeline="allow-pattern-rollback=0" -debug-only=serialize-to-isa \
 // RUN:  2>&1 | FileCheck %s
 
-// CHECK-LABEL: Generated by LLVM NVPTX Back-End
+// CHECK-LABEL: {{Generated by (LLVM NVPTX Back-End|NVIDIA NVVM Compiler)}}
 // CHECK: .visible .func kernel_a()
 // CHECK: ret;
 gpu.module @bar {
@@ -12,7 +12,7 @@ gpu.module @bar {
   }
 }
 
-// CHECK-LABEL: Generated by LLVM NVPTX Back-End
+// CHECK-LABEL: {{Generated by (LLVM NVPTX Back-End|NVIDIA NVVM Compiler)}}
 // CHECK: .visible .func  ({{.+}}) fma(
 // CHECK: fma.rn.f32
 

diff  --git a/mlir/test/Integration/GPU/CUDA/sm90/transform-dialect/tma_load_64x8_8x128_noswizzle-transform.mlir b/mlir/test/Integration/GPU/CUDA/sm90/transform-dialect/tma_load_64x8_8x128_noswizzle-transform.mlir
index a1e2729146c64..6ba9c16390192 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/transform-dialect/tma_load_64x8_8x128_noswizzle-transform.mlir
+++ b/mlir/test/Integration/GPU/CUDA/sm90/transform-dialect/tma_load_64x8_8x128_noswizzle-transform.mlir
@@ -19,10 +19,10 @@
 
 // Basic PTX check to make sure we are generating the right instructions.
 // CHECK-PTX: mbarrier.init.shared.b64
-// CHECK-PTX: mbarrier.arrive.expect_tx.shared.b64
-// CHECK-PTX: cp.async.bulk.tensor.2d.shared::cluster.global.tile.mbarrier::complete_tx::bytes
-// CHECK-PTX: cp.async.bulk.tensor.2d.shared::cluster.global.tile.mbarrier::complete_tx::bytes
-// CHECK-PTX: mbarrier.arrive.expect_tx.shared.b64
+// CHECK-PTX-DAG: mbarrier.arrive.expect_tx.shared.b64
+// CHECK-PTX-DAG: cp.async.bulk.tensor.2d.shared::cluster.global.tile.mbarrier::complete_tx::bytes
+// CHECK-PTX-DAG: cp.async.bulk.tensor.2d.shared::cluster.global.tile.mbarrier::complete_tx::bytes
+// CHECK-PTX-DAG: mbarrier.arrive.expect_tx.shared.b64
 // CHECK-PTX: mbarrier.try_wait.parity.shared.b64
 
 // RUN: mlir-opt %s \