在某些情况下,反汇编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))with i = tid - 6,以便现在可以使用所涉及的指令
ISETP.LT.U32.AND P0, PT, R0, 0x4, PT;
@P0如果谓词寄存器0为真,则在指令有条件地执行指令之前.同样,@!P0在指令之前意味着如果谓词寄存器0为假,则有条件地执行该指令.在反汇编更复杂的机器代码时,您将看到通常使用多个谓词寄存器.该预测机制还通过预测BRA指令用于条件分支.
ISETP是一个整数比较(此处:LT=小于),结果写入谓词寄存器.它允许链接对复合分支有用的谓词.在您的示例中,未使用链接,因为编译器使用了一个巧妙的转换,允许使用单个值来评估复合条件ISETP.这里,链接运算符是AND,并且生成的谓词ISETP与PT(= true)链接.我不确定第二个实例的重要性是什么PT,你可以通过检查其他用法示例找到答案.
PSETP与类似的行一样工作ISETP,但是对谓词而不是整数起作用.我没有必要仔细查看这条指令,因为它似乎不经常发生.根据我的判断,PSETP结合两个谓词寄存器并将结果存储到谓词寄存器.这里它结合!P0和PT经由(=真)AND.看来这条指令也支持链接,在这种情况下使用AND链接PT.您的例子代表逻辑否定的解释P0 = !P0似乎是正确的.在这种情况下ISETP,我不确定第三个意思是什么PT.