CUDA装配说明

Question

解压缩CUDA代码似乎是一个非常有用的工具，即使不是在某些情况下唯一的工具，以了解编译器的行为以及性能指标。

我要说的是，不幸的是，通过CUDA二进制实用工具Application Note提供的文档没有为用户提供解释CUDA组装说明所需的所有工具，或者至少我无法从该文档中推断出所有需要的信息。“数据自动化系统手册”没有提供比“数据自动化系统二进制实用指南”更多的信息。例如，我应该如何解释这些指令？

ISETP.LT.AND P0, PT, R3, RZ, PT;

和

PSETP.AND.AND P0, PT, !P0, PT, PT;

@P0在指令之前做什么？它是一个指令标签吗?如果谓词寄存器P0是真的，那么执行会跳到该标签上吗？有什么通用的方法来解释数据自动化系统的组装说明吗？

非常感谢。

编辑NJUFFA的评论

我编译了以下简单的内核

__global__ void test_kernel(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if ((tid > 5) & (tid < 10)) a[tid] = tid;
    else b[tid] = tid;
}

这导致

/*0000*/        MOV R1, c[0x1][0x100];                 /* 0x2800440400005de4 */
/*0008*/        S2R R0, SR_CTAID.X;                    /* 0x2c00000094001c04 */
/*0010*/        S2R R2, SR_TID.X;                      /* 0x2c00000084009c04 */
/*0018*/        IMAD R2, R0, c[0x0][0x8], R2;          /* 0x2004400020009ca3 */
/*0020*/        IADD R0, R2, -0x6;                     /* 0x4800ffffe8201c03 */
/*0028*/        ISETP.LT.U32.AND P0, PT, R0, 0x4, PT;  /* 0x188ec0001001dc03 */
/*0030*/        I2F.F32.S32 R0, R2;                    /* 0x1800000009201e04 */
/*0038*/   @!P0 ISCADD R3, R2, c[0x0][0x24], 0x2;      /* 0x400040009020e043 */
/*0040*/    @P0 ISCADD R2, R2, c[0x0][0x20], 0x2;      /* 0x4000400080208043 */
/*0048*/   @!P0 ST [R3], R0;                           /* 0x9000000000302085 */
/*0050*/    @P0 ST [R2], R0;                           /* 0x9000000000200085 */
/*0058*/        EXIT ;                                 /* 0x8000000000001de7 */

编译器已将条件((tid > 5) & (tid < 10))重铸为((i < 4) & (i >= 0))，并使用i = tid - 6，以便现在所涉及的指令是

ISETP.LT.U32.AND P0, PT, R0, 0x4, PT;

Answer 1

如果谓词寄存器0为真，则在指令有条件地执行该指令之前，@P0。同样，如果谓词寄存器0为false，则指令之前的@!P0意味着有条件地执行该指令。在分解更复杂的机器代码时，您将看到通常有多个谓词寄存器在使用。该预测机制还用于条件分支，通过预测BRA指令。

ISETP是一个整数比较(此处：LT =小于)，其结果写入谓词寄存器。它允许对谓词进行链接，这对于复合分支非常有用。在您的示例中，没有使用链接，因为编译器使用了一个巧妙的转换，它允许使用单个ISETP计算复合条件。这里，链式操作符是AND，由ISETP生成的谓词用PT (= true)链接。我不确定PT的第二个实例的意义是什么，您可能可以通过查看其他用法示例来找出。

PSETP的工作方式类似于ISETP，但工作在谓词上而不是整数上。我没有必要仔细研究这一指示，因为它似乎并不经常发生。据我所知，PSETP组合了两个谓词寄存器，并将结果存储到谓词寄存器中。在这里，它通过!P0和PT (= true)通过AND组合。这条指令似乎也支持链接，在本例中使用AND与PT链接。您对示例表示逻辑否定P0 = !P0的解释似乎是正确的。与ISETP的情况一样，我不确定第三个PT的含义是什么。