)]}'
{
  "commit": "442c605b3f85b0728389abb83ad842db38d5108e",
  "tree": "67819605df3f6ebfb44018e7835f5517b4eaa9b2",
  "parents": [
    "df5f71837e57f82435daa0a8678bdffb9d67bdac"
  ],
  "author": {
    "name": "WMC",
    "email": "tnwilly@gmail.com",
    "time": "Thu Jul 30 17:33:58 2026 +0800"
  },
  "committer": {
    "name": "GitHub",
    "email": "noreply@github.com",
    "time": "Thu Jul 30 11:33:58 2026 +0200"
  },
  "message": "[LLVMGPU] Add NVIDIA FP8 mma.sync intrinsics for sm_89 and sm_120 (#24659)\n\nPart of #24654\n\n## Summary\nAdds NVIDIA FP8 `mma.sync` (m16n8k32) strategy-selection support for Ada\n(sm_89) and Blackwell (sm_120). IREE previously had no NVIDIA FP8\nsupport at all — FP8 matmuls silently fell back to SIMT tiling on both\ntargets.\n\n## Not included\nFP8 lowering pipeline support. The TileAndFuse LHS transpose is\nhardcoded for element[K]\u003d2 (F16); FP8\u0027s element[K]\u003d4 causes a\nvector.shape_cast shape mismatch. Will do this in the next PR.\n\n## Test (on RTX PRO 6000 Blackwell, sm_120)\n\n```\nPATH\u003d/workspace/iree-build/tools:/workspace/iree-build/llvm-project/bin:$PATH \\\n  ./llvm-project/bin/llvm-lit -v \\\n  /workspace/compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm89.mlir \\\n  /workspace/compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir \\\n  /workspace/compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir\n```\n\n```\n-- Testing: 3 tests, 3 workers --\nPASS: IREE :: src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm89.mlir (1 of 3)\nPASS: IREE :: src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir (2 of 3)\nPASS: IREE :: src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir (3 of 3)\n\nTesting Time: 0.20s\n\nTotal Discovered Tests: 3\n  Passed: 3 (100.00%)\n```\n\n---------\n\nSigned-off-by: weimin023 \u003ctnwilly@gmail.com\u003e",
  "tree_diff": [
    {
      "type": "modify",
      "old_id": "158affdd4b10f11a0ef8a5a69c3694e289c746a4",
      "old_mode": 33188,
      "old_path": "compiler/plugins/target/CUDA/test/target_device_features.mlir",
      "new_id": "2fb107214aa32d6c87fbfa38d268dc303928124c",
      "new_mode": 33188,
      "new_path": "compiler/plugins/target/CUDA/test/target_device_features.mlir"
    },
    {
      "type": "modify",
      "old_id": "cc8f1dc302b7befdb485dc13140edc50354ab739",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp",
      "new_id": "0bc2968f3b94458c7fb9dcecb9bd8c325eb15ef4",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp"
    },
    {
      "type": "modify",
      "old_id": "19a64bfa3b1633e866750f89b0265203f59f72b5",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td",
      "new_id": "80f631139132fe533354cff617a72ed15c391926",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUEnums.td"
    },
    {
      "type": "modify",
      "old_id": "39724bfef41194f290b4c716f2532fcef6a10f4f",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp",
      "new_id": "65edff20e4369b9179c86445336c61efcff0eba5",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.cpp"
    },
    {
      "type": "modify",
      "old_id": "ae991c26bf861ea97f600f739afaafdd26cb402c",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/BUILD.bazel",
      "new_id": "5bc165f46ccc3216b198f1ae74dcacbcefb8939b",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/BUILD.bazel"
    },
    {
      "type": "modify",
      "old_id": "b293f5ddda8dbfef5ce562a0ad31f0512166bafd",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/CMakeLists.txt",
      "new_id": "6ca9b6d0568227c2c9b0e333ff99976484e35bba",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/CMakeLists.txt"
    },
    {
      "type": "modify",
      "old_id": "684c23a2a5c42dd8b5dcdd32ef5df37ed01fa11e",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir",
      "new_id": "aa6fed130feb387beebe2c1eef24980fc1a99cfc",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm120.mlir"
    },
    {
      "type": "add",
      "old_id": "0000000000000000000000000000000000000000",
      "old_mode": 0,
      "old_path": "/dev/null",
      "new_id": "4b4fd0972188f3bf365750c3f8a4e06bc514546a",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/config_tile_and_fuse_sm89.mlir"
    }
  ]
}