9.7.16.3. Major-ness supported by Strides

There are two strides involved while accessing a matrix from shared memory:

• Leading dimension stride (byte offset or absolute address)
• Stride dimension byte offset

9.7.16.3.1. Leading Dimension Stride: relative offset or absolute address

There are two modes of Leading Dimension Strides as described below. Bit #52 in the Shared memory descriptor is used to distinguish between two modes.

9.7.16.3.1.1. Relative offset mode

In this mode, the leading dimension stride is specified as a relative byte offset between the columns as explained in the below table. The leading dimension stride can either be specified as a relative offset between the columns or as an absolute byte address of next buffer. The leading dimension stride is defined differently for transposed and non-transposed matrices. The leading dimension stride is defined as follows for matrices whose element types are normalized to 128-bits:

9.7.16.3.1.2. Absolute address mode for K dimension being 48B

The tcgen05.mma instruction with K-dimension of 48B would overflow the 128B shared memory boundary if the data is packed contiguously.

In this case, the absolute address mode can be used to break up the data in the shared memory into two chunks such that both these chunks are laid out within the aligned 128-byte address boundary. The leading dimension absolute address can point to the second data chunk in the shared memory.

9.7.16.3.1.2.1. Restrictions on the Leading Dimension Absolute Address Stride

Following are the restrictions on the absolute address stride mode:

• Only 128B swizzle (with 16B atomicity) is supported.
• Only K-Major mode is supported. That is, the transpose bits (bits #15 and #16) in Instruction descriptor must be 0.
• The matrix base offset must be 0.

9.7.16.3.2. Stride Dimension Byte Offset

The stride dimension byte offset is defined differently for transposed and non-transposed matrices. The stride dimension byte offset is defined as follows for matrices whose element types are normalized to 128-bits:

9.7.16.3.3. Canonical Layouts

In terms of CuTe layouts the canonical layout can be expressed as follows:

where:

• T = 128 / sizeof-elements-in-bits — T represents scale factor which normalizes matrix element types to 128-bits.
• m represents the number of repeating patterns across rows.
• k represents the number of repeating patterns across columns.

Examples

K-Major, no-swizzling and tf32 type (Figure 188)

!K major, no-swizzling and tf32 type

Figure 188 K major, no-swizzling and tf32 type

The strides and related details are as follows:

• Exact layout: Swizzle<0,4,3> o ((8,2),(4,4)):((4,32),(1,64))
• Canonical Layout: Swizzle<0,4,3> o ((8,m),(T,2k)):((1T,SBO),(1,LBO))

K-Major, 32B swizzling and tf32 type (Figure 189)

!K major, 32B swizzling and tf32 type

Figure 189 K major, 32B swizzling and tf32 type

The strides and related details are as follows:

• Exact layout: Swizzle<1,4,3> o ((8,2),(4,4)):((8,64),(1,4))
• Canonical Layout: Swizzle<1,4,3> o ((8,m),(T,2k)):((2T,SBO),(1,T))

MN-Major, no-swizzling and bf16 type (Figure 190)

!MN major, no-swizzling and bf16 type

Figure 190 MN major, no-swizzling and bf16 type

The strides and related details are as follows:

• Exact layout: Swizzle<0,4,3> o ((8,1,2),(8,2)):((1,8,64),(8,128))
• Canonical Layout: Swizzle<0,4,3> o ((T,1,m),(8,k)):((1,T,SBO),(1T,LBO))

MN-Major, 32B swizzling and bf16 type (Figure 191)

!MN major, 32B swizzling and bf16 type

Figure 191 MN major, 32B swizzling and bf16 type

The strides and related details are as follows:

• Exact layout: Swizzle<1,4,3> o ((8,2,2),(8,2)):((1,8,128),(16,256))
• Canonical Layout: Swizzle<1,4,3> o ((T,2,m),(8,k)):((1,T,LBO),(2T,SBO))

MN-Major, 64B swizzling and bf16 type (Figure 192)

!MN major, 64B swizzling and bf16 type

Figure 192 MN major, 64B swizzling and bf16 type

The strides and related details are as follows:

• Exact layout: Swizzle<2,4,3> o ((8,4,2),(8,2)):((1,8,256),(32,512))
• Canonical Layout: Swizzle<2,4,3> o ((T,4,m),(8,k)):((1,T,LBO),(4T,SBO))