triton之内联汇编

一 输入输出

    constraints=(
        "=r,=r,=r,=r,=r,=r,=r,=r,"
        "r,r,r,r,r"
    ),

constraints：约束字符串，定义输入输出寄存器的分配。

• =r,=r,=r,=r,=r,=r,=r,=r,：输出寄存器，即 $0 到 $7。
• r,r,r,r,r：输入寄存器，即 $8 到 $12。

二 数据解包

{
    // Unpack `a` into `ai`.
    .reg .b8 tmp<4>;
    mov.b32 {tmp0, tmp1, tmp2, tmp3}, $8;
    cvt.u32.u8 $0, tmp0;
    cvt.u32.u8 $1, tmp1;
    cvt.u32.u8 $2, tmp2;
    cvt.u32.u8 $3, tmp3;
}

• 这里使用了 PTX 汇编语言，首先将 a 解包成四个 8 位寄存器。
• mov.b32 {tmp0, tmp1, tmp2, tmp3}, $8;：将 a 的内容移动到 4 个 8 位寄存器 tmp0 到 tmp3。
• cvt.u32.u8：将每个 8 位值转换为 32 位无符号整数。

每个寄存器只能装一个byte的数据