tree 2184ecf847f3902e0b389ccd909d5216de49dfa3
parent b32ca40857b9e91ed99f91ba86a730a1feb9ff49
author WMC <tnwilly@gmail.com> 1780986950 +0800
committer GitHub <noreply@github.com> 1780986950 +0200
gpgsig -----BEGIN PGP SIGNATURE-----
 
 wsFcBAABCAAQBQJqJ7RGCRC1aQ7uu5UhlAAAP6EQAHmiWgK2JUjRRFz1sYO5uKJj
 XpO89XxPdx76X5WhuzXNiwuMlZBSiIgzJiFnXwBFHXovOFgd6t9UUp9aHZCS5ojZ
 M4kaMeaedZBtyAuBQAFjWiWPH7IreS/dccmzQuOsyryK7Ja2GN8ntplPS8nqlvdT
 mMMWyPJnzulvk2jWYlmoDfUsBUgJUVC87p01qZWtCf0dYndRUHTaUmcZ3El/dRuS
 Q7M8eaFTMKOJN2Iq+PFwhiCXBcm4AobuDm34mzkaVfyetk6doEVVY8WA2B+9o4fS
 o/4mH/bx8fnbOaA72RRh2jyCsLre8pMWoH210yiEzZ2HN1QqPRtJt7+c/g+a0ERW
 IexhJqXX0mtbTQ6B4eKGl6jeoNSRNCuZwUtdhk/GJRwa+9khKf1vW9i1wV3zK13Z
 Na9S9jmP9y1Uic3JEYP4VjAGAAqzfZ0Knh1Jms3o8ZEKM1B9RGiC81hMuTF052w4
 /Eh5oAhBCmkct9SAC8pd28FFmFsYAqxQtpfBaStKqV09AhmhoBVxoKHzkituTt8S
 uM2ivaC9E3BuXt/RVIVSiFQLkbRHfgJODstQZSjS5tCtPyzE1kh/Bs8SIOdfnUcc
 tfMxVSpxoITWKuGnlIjemm/Frn8tc6X02PXEzLB0BAc3IWozmvjKTR2yCQNOgdFP
 kA2TKZ6mSmWjm7rXBMek
 =HhqR
 -----END PGP SIGNATURE-----
 

[LLVMGPU] Add sm_120 TileAndFuse config coverage, consolidate MMA sync pipeline tests (#24577)

## Description
Adds LLVMGPU TileAndFuse config-selection coverage for the CUDA `sm_120`
target and consolidates MMA sync pipeline-lowering coverage across
NVIDIA
GPU architectures.

The `sm_120` config-selection test verifies that f16 matmul cases select
the
NVIDIA MMA TileAndFuse path before any architecture-specific tuning is
introduced. The pipeline-lowering tests are kept target-independent
because
they start from IR with an attached `mma_sync` intrinsic and only verify
that
the selected lowering pipeline generates the corresponding
`nvgpu.mma.sync`
ops.

Covered `sm_120` config-selection cases:
- `f16xf16 -> f32` matmul selects TileAndFuse with `NV_MMA_SYNC`
- `f16xf16 -> f16` matmul selects TileAndFuse with `NV_MMA_SYNC`
- `f16xf16 -> f32` matmul accumulate sets `convert_acc_gemm`

Covered target-independent pipeline-lowering cases:
- `f16xf16 -> f32` TileAndFuse lowering generates `nvgpu.mma.sync`
- `f16xf16 -> f16` TileAndFuse lowering generates `nvgpu.mma.sync`
- `f16xf16 -> f32` VectorDistribute lowering generates `nvgpu.mma.sync`
- `f16xf16 -> f16` VectorDistribute lowering generates `nvgpu.mma.sync`

This is baseline regression coverage only; it does not introduce new
tile
heuristics or `sm_120`-specific tuning.

## Test
- `llvm-lit -v \

compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir
\

compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm80.mlir
\

compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir
\

compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_vector_distribute_mma_sync.mlir`


```
-- Testing: 4 tests, 4 workers --
PASS: IREE :: src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir (1 of 4)
PASS: IREE :: src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm80.mlir (2 of 4)
PASS: IREE :: src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_vector_distribute_mma_sync.mlir (3 of 4)
PASS: IREE :: src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir (4 of 4)

Testing Time: 0.14s

Total Discovered Tests: 4
  Passed: 4 (100.00%)
```

---------

Signed-off-by: weimin023 <tnwilly@gmail.com>