Hello, I was wondering about the .alayout and .blayout modifiers on mma instructions. The layout seems to be very strict and described in sections 9.7.13.4.[1-13] of the PTX ISA specification. Could I get some information on what would changing the layout would do ? I’ve only ever seen .row.col In…

I don’t think there is a repository of sample codes that use/demonstrate each and every variant of each PTX mma instruction. CUTLASS tends to use them to some degree, and with a bit of searching you can find a few examples here on these forums. The .row and .col refer to row-major and column-major …

Thanks for your answer! I guess what I’m mostly confused is that to me, row-major and column-major make sense to describe how a matrix is stored in a contiguous memory region. For those instructions, the inputs are expected to be in registers spread over the threads in the warp and which register co…

[image] leclerc1: For those instructions, the inputs are expected to be in registers spread over the threads in the warp and which register contain what part of the matrix is strictly specified in the sections above. For the wmma instructions I think the information makes sense. For the mma …

I understand that wmma doesn’t have a visible layout and I’m not talking about those here. I also understand why a load (ldmatrix operation) would need the .alayout and .blayout modifiers to properly place the elements in registers. What I don’t understand is why the multiplication part (mma opera…

[image] leclerc1: What should change in this code if I swaped row for col ? When I attempt to change either .row to .col or .col to .row in the complete code from here , I get a PTX error, basically saying that is an illegal choice or an illegal instruction. ptxas /tmp/tmpxft_00040ec9_0000000…

I’m trying to understand the specification itself, which mentions two modifiers .row and .col which behavior/implication isn’t well defined (and in the current implementation illegal to change as you seem to suggest), so I wouldn’t say the question is moot. If what you are saying is correct I think …

[image] leclerc1: I think the document would needs an update. Anyone can file a bug at any time to request document clarifications.

Exchanging row and col would access the data and combine matrices in a matrix multiplication in a totally different way. That is not how the architecture of SMs are made. I for my part do not understand mma in terms of rows and cols, but in terms of dimension which is added over and (two) independe…

Confused about (a,b)layout on mma.sync instructions

Accelerated Computing CUDA CUDA Programming and Performance

Robert_Crovella March 20, 2024, 8:59pm 8

Anyone can file a bug at any time to request document clarifications.

Topic		Replies	Views
How to use mma instructions to mul matrix in col major CUDA Programming and Performance	5	444	November 23, 2024
Does wmma::load_matrix_sync() or wmma::mma_sync() performance benefit from row-major or col-major layouts? CUDA Programming and Performance cuda	1	453	November 15, 2025
Understand the mma instruction in PTX CUDA Programming and Performance	4	1510	May 29, 2024
What does the "row" and "col" mean in mma.sync.aligned.m16n8k4.row.col.f32.tf32.tf32.f32 CUDA Programming and Performance	0	303	October 30, 2023
Address out of bounds when using mma instructions CUDA Programming and Performance cuda	7	396	June 11, 2024
Questions about mma instruction with Nvidia ptx CUDA Programming and Performance cuda	1	313	July 15, 2024
Fastest Tiled WMMA for Matrices of Any Size? CUDA Programming and Performance	3	576	October 26, 2024
Clarification on the accumulator layout in an mma instruction CUDA Programming and Performance	1	423	November 30, 2023
Alignment requirement for the `ldmatrix` instruction CUDA Programming and Performance cuda , instruction	2	419	November 1, 2024
The fragment layout of multiplicand A is not clear in mma.sp.sync.aligned.m16n8k32 when data type is fp16/bf16 CUDA Programming and Performance	2	766	February 22, 2022

Confused about (a,b)layout on mma.sync instructions

Related topics