NVIDIA / NVIDIA/cutlass

[QST]Question about memory layout compatibility between mma.sync.m16n8k64.mxf4nvf4 and tcgen05 128×128×64 instructions

Open
#2,951 4 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage inactive-30d inactive-90d question
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

I have a question regarding memory layout compatibility between the following PTX instruction:

mma.sync.aligned.m16n8k64.row.col
.kind::mxf4nvf4.block_scale.scale_vec::4X
.f32.e2m1.e2m1.f32.ue4m3

and the newer tcgen05 128×128×64 Tensor Core instructions.

Main Question

On current GPU architectures(SM101), only tcgen05 Tensor Core instructions are supported, while the above mma.sync(SM120+) instruction is no longer available.

Therefore, my goal is to replace the existing mma.sync.m16n8k64 instruction sequence with tcgen05 instructions (for example, a single tcgen05 128×128×64 operation, or another supported mnk shape).

The key constraint is that I would like to understand whether this replacement can be done without changing the existing memory layout of A, B, and their block scales.

Is it possible to directly replace mma.sync.aligned.m16n8k64.mxf4nvf4.block_scale.scale_vec::4X with a tcgen05 instruction (e.g. 128×128×64 or another shape), while reusing the same memory layout for:

  • A operand fragments
  • B operand fragments
  • Block-scale vectors and their indexing

In other words:

  • Does tcgen05 assume a compatible fragment ordering, row/col layout, and scale grouping as the m16n8k64 MMA instruction?

  • Or does tcgen05 require a different packing / swizzling / scale layout, making a full repacking of A/B matrices and scale vectors unavoidable?

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start by comparing the documented fragment ordering, row/column layouts, packing, swizzling, and scale grouping for mma.sync.aligned.m16n8k64.mxf4nvf4 and the tcgen05 128×128×64 instructions. The issue names no repository files or tests, so the research must first locate the relevant PTX and CUTLASS instruction documentation. Done means a maintainer-confirmed compatibility answer covering A, B, and block-scale layouts.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc
Issue type
Documentation
Difficulty
5/5
Estimated time
Over a week
Activity status
Quiet
Clarity
Needs clarification
Newbie friendliness
25/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.