Вопросы по инструкции mma с Nvidia ptx

Привет, насколько я понимаю инструкцию MMA с ptx (пожалуйста, скажите мне, если я ошибаюсь):

1. это инструкция для варпа
2. необходимо загрузить определенный элемент в регистр каждого потока внутри целевой деформации.
3. Размер умножения фиксирован (ограниченный набор на выбор)

У меня есть скомпилированный набор инструкций ptx, и я выбрал одну инструкцию mma:

//
// Generated by LLVM NVPTX Back-End
//

.version 8.4
.target sm_89
.address_size 64

    // .globl   matmul_kernel
.extern .shared .align 16 .b8 global_smem[];

.visible .entry matmul_kernel(
    .param .u64 matmul_kernel_param_0,
    .param .u64 matmul_kernel_param_1,
    .param .u64 matmul_kernel_param_2,
    .param .u32 matmul_kernel_param_3,
    .param .u32 matmul_kernel_param_4,
    .param .u32 matmul_kernel_param_5,
    .param .u32 matmul_kernel_param_6,
    .param .u32 matmul_kernel_param_7,
    .param .u32 matmul_kernel_param_8
)
.maxntid 128, 1, 1
{
    ...
    ldmatrix.sync.aligned.m8n8.x4.shared.b16 { %r3100, %r3101, %r3102, %r3103 }, [ %r561 + 0 ];
    ...
    ldmatrix.sync.aligned.m8n8.x4.trans.shared.b16 { %r3084, %r3085, %r3086, %r3087 }, [ %r581 + 0 ];
    ...
    mov.f32     %f2306, 0f00000000;
    mov.b32     %r3107, 2;
    mov.b32     %r3106, 0;
    shl.b32     %r2885, %r100, 1;
    shl.b32     %r2894, %r101, 1;
    shl.b32     %r2895, %r102, 1;
    shl.b32     %r2896, %r103, 1;
    shl.b32     %r2897, %r104, 1;
    shl.b32     %r2898, %r105, 1;
    shl.b32     %r2899, %r106, 1;
    mov.u32     %r3104, %r765;
    mov.u32     %r3105, %r758;
    mov.f32     %f2307, %f2306;
    mov.f32     %f2308, %f2306;
    mov.f32     %f2309, %f2306;
    ...
    mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 { %f2306, %f2307, %f2308, %f2309 }, { %r3100, %r3101, %r3102, %r3103 }, { %r3084, %r3085 }, { %f2306, %f2307, %f2308, %f2309 };
}

Теперь мои вопросы:

1. по инструкции это умножение D{16, 8}=A{16, 16}B{16, 8}+D{16, 8}$ с A,B в типе f16 и D в типе f32. Но фактическая инструкция загрузки загружает только 4 регистра для A. Означает ли это, что для каждого потока в warp она загружает только 8 элементов (поскольку ширина регистра составляет 32 бита), а всего 32 потока будут загружать 8x32 = 16x16 элементов A? и то же самое для матрицы B и матрицы результатов D (D — 4x32=16x8 в типе f32)
2. означает ли тег sync, что целевая деформация будет продолжаться до тех пор, пока матрица не завершится тензорным ядром?
3. видно, что версия ptx — 8.4, а из официального документа видно, что 8.4 поддерживает только форму m16n8k64 для разреженной и .m16n8k32 для плотной формы, для m16n8k16 нет поддержки, или m16n8k64 означает, что она поддерживает любую форму, которая m<16 и n<8 и k<64?
4. как упоминалось выше, что означает плотная и разреженная матрица?