)]}'
{
  "commit": "e97b2408196bc5bac0c590ed25eff2f65af966d6",
  "tree": "17b9173fb3e5fc27c9d7f3e9aa8b2947bf275ee4",
  "parents": [
    "6ddc991778bf91df647bec3946356428d9109109"
  ],
  "author": {
    "name": "WMC",
    "email": "tnwilly@gmail.com",
    "time": "Thu Oct 08 16:36:11 2026 +0800"
  },
  "committer": {
    "name": "GitHub",
    "email": "noreply@github.com",
    "time": "Thu Oct 08 10:36:11 2026 +0200"
  },
  "message": "[LLVMGPU] Fix hardcoded LHS transpose shape in NV mma.sync lowering (#24676)\n\nThis PR is stacked on:\n- [x] https://github.com/iree-org/iree/pull/24659\n\n## Description\n\n`createMmaOp()` hardcoded the LHS reshape shape to `{2, 2, 2}` when\nlowering NV `mma.sync`. This held for all NV intrinsics that existed at\nthe time (F16/BF16 m16n8k16, `element[K]\u003d2`), but FP8 m16n8k32 packs\n`element[K]\u003d4`, so the reshape produced 16 elements while the hardcoded\ntarget expected 8:\n\n\n\u003e error: \u0027vector.shape_cast\u0027 op has different number of elements at\nsource (16) and result (8)\n\u003e note: see current operation: %73 \u003d \"vector.shape_cast\"(%70) :\n(vector\u003c2x1x2x4xf8E4M3FN\u003e) -\u003e vector\u003c2x2x2xf8E4M3FN\u003e\n\nThis computes the reshape shape from `getSingleSubgroupLayout()` instead\nof hardcoding it, so it adapts to any `element[K]` size. F16/BF16\nbehavior is unchanged (`{2, 2, 2}` still holds for them).\n\nAdds an FP8 (E4M3FN) TileAndFuse pipeline test.\n\n## Testing\n\n```bash\n/workspace/iree-build/tools/iree-opt \\\n    --split-input-file \\\n    --iree-codegen-llvmgpu-nvvm-lowering-pipeline\u003d\"include-llvm-lowering\u003dfalse\" \\\n    compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir \\\n    | /workspace/iree-build/llvm-project/bin/FileCheck \\\n      compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir\n\n---------\n\nSigned-off-by: weimin023 \u003ctnwilly@gmail.com\u003e",
  "tree_diff": [
    {
      "type": "modify",
      "old_id": "0bc2968f3b94458c7fb9dcecb9bd8c325eb15ef4",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp",
      "new_id": "d92c69bfa442ab10c84b1cda0388efd1090b28c6",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/Dialect/GPU/IR/IREEGPUAttrs.cpp"
    },
    {
      "type": "modify",
      "old_id": "99fa509ad5e5e5319f337fe8502cf826bb687b7d",
      "old_mode": 33188,
      "old_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir",
      "new_id": "c32b70b0eb444e9ac71be02a096a9596d6d10fd0",
      "new_mode": 33188,
      "new_path": "compiler/src/iree/compiler/Codegen/LLVMGPU/test/NVVM/pipeline_tile_and_fuse_mma_sync.mlir"
    }
  ]
}