[all-commits] [llvm/llvm-project] 4362da: [AMDGPU] Add wait states between different MFMAs s...

[Pankaj Dwivedi via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 4362da25049882128f50a0f6b4fb69c1a0f8fd57
      https://github.com/llvm/llvm-project/commit/4362da25049882128f50a0f6b4fb69c1a0f8fd57
  Author: Pankaj Dwivedi <pankajkumar.divedi at amd.com>
  Date:   2026-09-28 (Mon, 28 Sep 2026)

  Changed paths:
    M llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
    M llvm/lib/Target/AMDGPU/GCNHazardRecognizer.h
    M llvm/test/CodeGen/AMDGPU/mai-hazards-gfx90a.mir
    M llvm/test/CodeGen/AMDGPU/mai-hazards-gfx942.mir

  Log Message:
  -----------
  [AMDGPU] Add wait states between different MFMAs sharing an accumulator (#218363)

A full-register src2/C read after an MFMA write emits no wait states,
relying on accumulator forwarding that only works while the chain stays
on one MFMA. Two different MFMAs sharing an accumulator instead need the
wait states of a partial overlap: on gfx950, v_mfma_f32_16x16x32_f16
then
v_mfma_f32_16x16x16_f16 on the same tuple needs 5 and got none. Compare
canonicalized opcodes so mac and register-bank forms still match.

Assisted-by: Claude Opus 5



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications