Какова именно раскладка значений потоков cute? ⇐ C++

[Anonymous]

Я изучаю структуру значений потоков Cute и следую блогу Леймао:
https://leimao.github.io/blog/CuTe-Thread-Value-Layout/
Я хочу выяснить, какой поток какую часть данных в матрице читает.
Итак, давайте возьмем SM80_16x8x16_F16F16F16F16_TN, например
когда вы печатаете этот атом мма, вы получаете следующую информацию:

Код: Выделить всё

    MMA_Atom
ThrID:      _32:_1
Shape_MNK:  (_16,_8,_16)
LayoutA_TV: ((_4,_8),(_2,_2,_2)):((_32,_1),(_16,_8,_128))
LayoutB_TV: ((_4,_8),(_2,_2)):((_16,_1),(_8,_64))
LayoutC_TV: ((_4,_8),(_2,_2)):((_32,_1),(_16,_8))

когда вы print_latex(mma), вы получаете это латексное изображение:


чья матрица A (16x16) отображается в левой части.
Итак, я печатаю LayoutA_TV, используя этот код:

Код: Выделить всё

auto s2xh4 = cute::make_layout(cute::make_shape (cute::make_shape (4,8),cute::make_shape (2,2,2)),cute::make_stride(cute::make_stride(32,1),cute::make_stride(16,8,128)));
cute::print_latex(s2xh4);

и я получил вот такой макет


Это означает:
для потока 0 он прочитал эти 8 значений:
0,16,8,24,128,144,136,152
если A является основным столбцом в структуре памяти, эти 8 индексов вычисляют координату следующим образом:
(0,0) (0,1) (8,0),(8,1),(0,8) (0,9),(8,8),(8,9)
это именно то, что отображается в мма изображение :
T0V0 в (0,0) T0V1 в (0,1) .... T0V4 в (0,8) T0V5 в (0,9).....
но в соответствии с определениями SM80_16x8x16_F16F16F16F16_TN:

Код: Выделить всё

struct SM80_16x8x16_F16F16F16F16_TN
{
using DRegisters = uint32_t[2];
using ARegisters = uint32_t[4];
using BRegisters = uint32_t[2];
using CRegisters = uint32_t[2];

CUTE_HOST_DEVICE static void
fma(uint32_t      & d0, uint32_t      & d1,
uint32_t const& a0, uint32_t const& a1, uint32_t const& a2, uint32_t const& a3,
uint32_t const& b0, uint32_t const& b1,
uint32_t const& c0, uint32_t const& c1)
{
#if defined(CUTE_ARCH_MMA_SM80_ENABLED)
asm volatile(
"mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16 "
"{%0,  %1},"
"{%2,  %3,  %4,  %5},"
"{%6,  %7},"
"{%8,  %9};\n"
: "=r"(d0), "=r"(d1)
:  "r"(a0),  "r"(a1),  "r"(a2),  "r"(a3),
"r"(b0),  "r"(b1),
"r"(c0),  "r"(c1));
#else
CUTE_INVALID_CONTROL_PATH("Attempting to use SM80_16x8x16_F16F16F16F16_TN without CUTE_ARCH_MMA_SM80_ENABLED");
#endif
}
};

это инструкция ptx — mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16
это означает, что матрица A является старшей строкой.
тогда это совершенно неправильно
Я знаю, что-то не так с моим пониманием.
Итак, где я неправильно понимаю?

Подробнее здесь: https://stackoverflow.com/questions/798 ... lue-layout