CUDA 组装说明

Posted 2023-02-16

【中文标题】CUDA 组装说明

【英文标题】：CUDA assembly instructions

【发布时间】：2013-10-14 09:44:22

【问题描述】：

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

我会说不幸的是，通过 CUDA BINARY UTILITIES Application Note 提供的文档并没有为用户提供解释 CUDA 汇编指令所需的所有工具，或者至少我无法从该文档中推断出所有需要的信息. “CUDA 手册”一书没有提供比 CUDA BINARY UTILITIES 指南更多的信息。例如，我应该如何解释说明

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

和

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

@P0 在指令前做了什么？如果谓词寄存器P0 为真，它是否是一个指令标签，以便执行跳转到该标签？有什么通用的方法来解释 CUDA 汇编指令吗？

非常感谢。

根据 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;

【问题讨论】：

@Px 和@!Px 装饰器（我认为）表示给定指令的条件执行。这不包括分支

@talonmies 谢谢。这也是我的理解。我同意没有跳转，只有条件执行。

我对第一条指令有一个猜测。也许PT 是一种分隔符。该指令总是根据结构（来自我在互联网上找到的）来，即ISETP.LT.AND P0, PT, Op1, Op2, PT;。在上述情况下，它只是简单地实现R3<RZ 并将比较结果存储在P0 中。我知道这样的指令应该实现R3<0，所以也许RZ 是一个包含0 的特殊寄存器。当然，这只是猜测。

我知道 PTX 和已分解的代码之间没有一对一的映射，但也许 PTX 可以帮助理解后者发生了什么。与上述第一条汇编指令对应的 PTX 指令是setp.lt.s32 %p1, %r1, 0;，它将r1<0 的结果放入谓词寄存器p1。

关于PSETP.AND.AND P0, PT, !P0, PT, PT; 指令，它显然将!P0 放入P0。对应的 PTX 指令是not.pred %p2, %p1;，应该将!p1 放入p2。其余的PT 是分隔符或一种用于填充操作数位置的中性命令。

【参考方案1】：

@P0 如果谓词寄存器 0 为真，则在指令有条件地执行指令之前。同样，指令前的@!P0 表示如果谓词寄存器0 为假，则有条件地执行该指令。在反汇编更复杂的机器代码时，您会看到通常使用多个谓词寄存器。这种预测机制也用于条件分支，通过预测BRA 指令。

ISETP 是一个整数比较（这里：LT = 小于），结果写入谓词寄存器。它允许链接谓词，这对复合分支很有用。在您的示例中，没有使用链接，因为编译器使用了一种巧妙的转换，允许使用单个 ISETP 评估复合条件。在这里，链接运算符为AND，ISETP 生成的谓词与PT 链接（= true）。我不确定PT 的第二个实例的意义是什么，您可以通过查看其他用法示例来了解。

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