Skip to content

fix: remove force nt path in tensorwise grouped gemm - #437

Open
RuibinCheung wants to merge 1 commit into
mainfrom
fix/zhangrb/remove_force_nt_path_in_gg
Open

fix: remove force nt path in tensorwise grouped gemm#437
RuibinCheung wants to merge 1 commit into
mainfrom
fix/zhangrb/remove_force_nt_path_in_gg

Conversation

@RuibinCheung

Copy link
Copy Markdown
Collaborator

Description

The non-NT layout grouped GEMM backends (triton and flydsl) are now performant enough on MI350 , so this workaround is no longer needed. This PR removes it and lets TENSORWISE grouped GEMM take the original code path on all architectures: the backward grad_a is computed directly from the untransposed weight with trans_b=not ctx.trans_b, matching what gfx942 has been doing all along.

This is a pure code-path simplification for the TENSORWISE granularity. The ROWWISE, BLOCKWISE and MX_BLOCKWISE paths are untouched, and the identically named helper in gemm_fp8.py (non-grouped GEMM) is intentionally left in place. Keeping non-NT layout in gemm because we still use hipblaslt backend in gemm. And the non-NT layout performance is not good enough in hipbalslt backend.

Fixes # (issue)

Type of change

  • Documentation change (change only to the documentation, either a fix or a new content)
  • Bug fix (non-breaking change which fixes an issue)
  • New feature (non-breaking change which adds functionality)
  • Breaking change (fix or feature that would cause existing functionality to not work as expected)
  • Infra/Build change
  • Code refactoring

Changes

Please list the changes introduced in this PR:

  • Removed the _deter_use_nt_layout_gemm_in_bwd helper from primus_turbo/pytorch/ops/grouped_gemm_fp8.py, together with the now-unused is_gfx942 import.
  • FP8GroupedGemmTensorFunc.forward: dropped the transposed-weight cache branch (both the b_t reuse and the transpose(-1, -2).contiguous() fallback), and now always saves the untransposed quantized_b for backward. The b_t argument is unused and annotated as such, consistent with a_t.
  • FP8GroupedGemmTensorFunc.forward: no longer stores use_nt_layout_gemm_in_bwd on ctx.
  • FP8GroupedGemmTensorFunc.backward: collapsed the two grad_a branches into the single original call with trans_b=not ctx.trans_b. The grad_b (variable-K wgrad) call is unchanged.

Checklist:

  • The functionality is complete
  • I have commented my code, particularly in hard-to-understand areas
  • I have made corresponding changes to the documentation
  • My changes generate no new warnings
  • I have added tests that prove my fix is effective or that my feature works
  • New and existing unit tests pass locally with my changes

Copilot AI review requested due to automatic review settings July 29, 2026 09:44

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

This PR simplifies the FP8 grouped GEMM (TENSORWISE) autograd path by removing the MI350/MI450-era backward workaround that forced an NT-layout code path. It standardizes backward grad_a computation to always use the original call form (trans_b=not ctx.trans_b) and always saves the untransposed quantized weight for backward.

Changes:

  • Removed the architecture-gated _deter_use_nt_layout_gemm_in_bwd helper and the related gfx942 check/import.
  • Simplified FP8GroupedGemmTensorFunc.forward to always save quantized_b (dropping the transposed-weight cache branch).
  • Simplified FP8GroupedGemmTensorFunc.backward to a single grad_a call path using trans_b=not ctx.trans_b.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants