)]}'
{
  "commit": "a4d0287253fbe7211f545cda304b487218c2a711",
  "tree": "ed15bd28cc2de4925e76a46dd3e72d6fb7c0bd1f",
  "parents": [
    "540008c4aa3f6a2c6e55a054a33f24bd85144d97"
  ],
  "author": {
    "name": "WMC",
    "email": "tnwilly@gmail.com",
    "time": "Tue Jun 30 19:35:33 2026 +0800"
  },
  "committer": {
    "name": "GitHub",
    "email": "noreply@github.com",
    "time": "Tue Jun 30 13:35:33 2026 +0200"
  },
  "message": "[CUDA][LLVMGPU] Fix sm_120 WGP params and add BF16 mma.sync coverage (#24648)\n\n## Description\n\nFixes a correctness issue in `getSM120WgpDetails()` carried over from\nthe Ampere placeholder, consolidates BF16 test coverage into existing\narchitecture-scoped files, and extends sm_120 tests with a\nmixed-precision case.\n\n- **BF16 mma.sync fix**: `NV_MMA_SYNC_F32_16x8x16_BF16` was missing from\nsm_120\u0027s `mmaOps` despite being supported from Ampere onwards. sm_121\nalready advertises this intrinsic. Without this fix, bf16 × bf16 → f32\nmatmuls silently fall back to SIMT tiling on sm_120.\n\n- **Test consolidation**: Per the reviewer\u0027s suggestion, tests no longer\nhave their own per-architecture files. Instead they are folded into the\nexisting architecture-scoped test files.\n\n## Testcases\n\n1. config_tile_and_fuse_sm80.mlir / config_tile_and_fuse_sm120.mlir:\neach gets a new BF16 section verifying that bf16 × bf16 → f32 selects\n`NV_MMA_SYNC_F32_16x8x16_BF16` at the config stage.\n\n1. config_tile_and_fuse_sm120.mlir: adds a mixed-precision f32 × bf16 →\nf32 case to verify the SIMT fallback path when no MMA intrinsic matches\nthe operand types.\n\n1. pipeline_tile_and_fuse_mma_sync.mlir: adds a shared BF16 pipeline\nlowering section that covers both sm_80 and sm_120, verifying\n`nvgpu.mma.sync` is generated with the correct shape.\n\n1. target_device_features.mlir: expands the SM120 metadata check to a\nfull wgp attribute check (mma list, workgroup sizes, memory limits),\nmatching the granularity of the existing SM89 check.\n\n1. bf16_mma_sm80.mlir: deleted — coverage moved into the files above.\n\n---------\n\nSigned-off-by: weimin023 \u003ctnwilly@gmail.com\u003e",
  "tree_diff": [
    {
      "type": "modify",
      "old_id": "91d4f40231e8f31a42a61a9a0cbd1d0131d8f2b5",
      "old_mode": 33188,
      "old_path": "compiler/plugins/target/CUDA/test/target_device_features.mlir",
      "new_id": "158affdd4b10f11a0ef8a5a69c3694e289c746a4",
      "new_mode": 33188,
      "new_path": "compiler/plugins/target/CUDA/test/target_device_features.mlir"
    },
    {
      "type": "modify",
      "old_id": "64ccd853f94c6b7c247da5b6018d5d1010b18a18",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp",
      "new_id": "39724bfef41194f290b4c716f2532fcef6a10f4f",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp"
    },
    {
      "type": "modify",
      "old_id": "e5ded3077e318915a91ab96a7afc334e9b8f06a0",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/BUILD.bazel",
      "new_id": "ae991c26bf861ea97f600f739afaafdd26cb402c",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/BUILD.bazel"
    },
    {
      "type": "modify",
      "old_id": "5d58dde3a93e8096b51a82d300a8057b2ae53658",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/CMakeLists.txt",
      "new_id": "b293f5ddda8dbfef5ce562a0ad31f0512166bafd",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/CMakeLists.txt"
    },
    {
      "type": "delete",
      "old_id": "6873631547c2263266723ae25866ef0b84708946",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/bf16_mma_sm80.mlir",
      "new_id": "0000000000000000000000000000000000000000",
      "new_mode": 0,
      "new_path": "/dev/null"
    },
    {
      "type": "modify",
      "old_id": "89cba205c09e2a206a496cdbe3488b3d36abf6e0",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir",
      "new_id": "684c23a2a5c42dd8b5dcdd32ef5df37ed01fa11e",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir"
    },
    {
      "type": "modify",
      "old_id": "69517d0c97ddea4e7c4875ec63838151d278f114",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm80.mlir",
      "new_id": "6bd2f94109ad833d1230397ce64d7742a2b3daf8",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm80.mlir"
    },
    {
      "type": "modify",
      "old_id": "5d99b252718ac75dd21602cd902d4ff100a4046c",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir",
      "new_id": "748597b270dbfe718ed66213fd769de995b822c4",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir"
    }
  ]
}