PTX 程序集加载/存储中寄存器周围括号的含义

Question

下面是由 Triton 编译器生成的明显合法的 PTX 汇编代码。我对加载和存储指令中使用的

{ %r1 }

和

{ %r2 }

感到困惑。根据 PTX ISA 文档，它看起来像一个初始化列表。但这没有意义。不仅因为初始化规范没有提到寄存器的使用。甚至不是因为加载/存储语义中的初始化器是无用的（没有什么可以初始化）。最重要的是，我对以下事实感到困惑：在加载/存储中使用

{}

会将参数的含义从标量更改为指针立即数。

也许，一个无聊的开发者只是想让大家的组装体验更加混乱。有没有人有更好的解释？

.version 7.5
.target sm_35
.address_size 64

        // .globl       E__01

.visible .entry E__01(
        .param .u64 E__01_param_0,
        .param .u64 E__01_param_1
)
.maxntid 128, 1, 1
{
        .reg .pred      %p<3>;
        .reg .b32       %r<4>;
        .reg .b64       %rd<3>;
        .loc    1 6 0
$L__func_begin0:
        .loc    1 6 0

        ld.param.u64    %rd2, [E__01_param_0];
        ld.param.u64    %rd1, [E__01_param_1];
        mov.pred        %p1, -1;
$L__tmp0:
        .loc    1 7 19
        mov.u32 %r1, 0x0;
        @%p1 ld.global.b32 { %r1 }, [ %rd1 + 0 ];
        .loc    1 8 18
        shl.b32         %r2, %r1, 1;
        .loc    1 9 22
        mov.u32         %r3, %tid.x;
        setp.eq.s32     %p2, %r3, 0;
        @%p2 st.global.b32 [ %rd2 + 0 ], { %r2 };
        .loc    1 9 2
        ret;
$L__tmp1:
$L__func_end0:

}

Answer 1

此语法似乎是加载向量和解包或打包向量和存储的有效概括。基本上它将单个寄存器视为大小为 1 的向量。

在这样的示例中可以看到具有适当向量的相同语法：

__global__ void foo(float2* arr, int n) {
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
    if (tid < n) {
        float2 x = arr[tid];
        arr[tid] = make_float2(x.x * x.y, x.x + x.y);
    }
}

其内部

if

编译为

        ld.global.v2.f32        {%f1, %f2}, [%rd4];
        mul.f32         %f3, %f1, %f2;
        add.f32         %f4, %f1, %f2;
        st.global.v2.f32        [%rd4], {%f3, %f4};