我已经使用asm volatile编写了以下一行CUDA PTX指令
__global__ void add( uint32_t a, uint32_t b )
{
uint32_t c = 0;
asm volatile("add.u32 %0, %1, %2;" : "=r"(c) : "r"(a), "r"(b) );
printf("sink=%d\n", c);
}
使用此内核调用
add<<< 1, 1 >>>( 1,2 );
和此编译命令
nvcc -arch=sm_70 -Xptxas -O3,-v,-dlcm=ca -o test1 test1.cu
我看到这个SASS代码
code for sm_70
Function : _Z6addjj
.headerflags @"EF_CUDA_SM70 EF_CUDA_PTX_SM(EF_CUDA_SM70)"
/*0000*/ IMAD.MOV.U32 R1, RZ, RZ, c[0x0][0x28] ; /* 0x00000a00ff017624 */
/* 0x000fd000078e00ff */
/*0010*/ @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ ; /* 0x000000fffffff389 */
/* 0x000fe200000e00ff */
/*0020*/ IMAD.MOV.U32 R0, RZ, RZ, c[0x0][0x160] ; /* 0x00005800ff007624 */
/* 0x000fe200078e00ff */
/*0030*/ IADD3 R1, R1, -0x8, RZ ; /* 0xfffffff801017810 */
/* 0x000fe40007ffe0ff */
/*0040*/ MOV R4, 0x0 ; /* 0x0000000000047802 */
/* 0x000fe40000000f00 */
/*0050*/ IADD3 R0, R0, c[0x0][0x164], RZ ; /* 0x0000590000007a10 */
/* 0x000fe40007ffe0ff */
/*0060*/ IADD3 R6, P0, R1, c[0x0][0x20], RZ ; /* 0x0000080001067a10 */
/* 0x000fe40007f1e0ff */
/*0070*/ MOV R5, 0x0 ; /* 0x0000000000057802 */
/* 0x000fc60000000f00 */
/*0080*/ IMAD.X R7, RZ, RZ, c[0x0][0x24], P0 ; /* 0x00000900ff077624 */
/* 0x000fe200000e06ff */
/*0090*/ STL [R1], R0 ; /* 0x0000000001007387 */
/* 0x0001ee0000100800 */
/*00a0*/ MOV R20, 0x0 ; /* 0x0000000000147802 */
/* 0x000fe40000000f00 */
/*00b0*/ MOV R21, 0x0 ; /* 0x0000000000157802 */
/* 0x000fd00000000f00 */
/*00c0*/ CALL.ABS.NOINC 0x0 ; /* 0x0000000000007943 */
/* 0x001fea0003c00000 */
/*00d0*/ EXIT ; /* 0x000000000000794d */
/* 0x000fea0003800000 */
/*00e0*/ BRA 0xe0; /* 0xfffffff000007947 */
/* 0x000fc0000383ffff */
/*00f0*/ NOP; /* 0x0000000000007918 */
/* 0x000fc00000000000 */
......................
我知道编译器会优化事情，但是，我希望看到2次加载，1次添加和1次存储。
我应该怎么做才能看到那个？另外，哪个SASS指令与我的add.u32指令完全对应？
CUDA内核调用涉及内核参数的按值传递。这些参数（作为内核调用机制的一部分）传递到设备，以供constant memory中的设备代码使用。
因此，需要任何a和b参数的设备代码都希望在常量存储器中找到它们。
我希望看到2次加载，1次添加和1次存储
粗略地说，这些都在那里。在某些情况下，它更易于向后工作，因为我们不知道编译器将使用哪个先验寄存器来注册代码中的哪些项目。
您的printf语句（一个函数调用）显然由该调用服务：
/*00c0*/ CALL.ABS.NOINC 0x0 ; /* 0x0000000000007943 */
该函数调用将需要某种“调用”框架。它将期望某种模式的堆栈，本地存储器和/或寄存器填充来完成其工作。我们在您的内核之前看到内核中唯一的本地存储：
/*0090*/ STL [R1], R0 ; /* 0x0000000001007387 */
因此，我们可以推测，加法运算的结果必须在寄存器R0中结束。我们将使用此信息来快速确定您要询问的添加指令的流程。那么我们的下一个问题是“哪个加法指令将结果存入R0？”那似乎就是这个：
/*0050*/ IADD3 R0, R0, c[0x0][0x164], RZ ; /* 0x0000590000007a10 */
（第一个操作数始终是目标操作数，内核中没有其他加法运算将R0作为目标）。我们知道此加法必须以某种方式将a和b加在一起。现在我们还知道a和b将被期望从常量存储器加载。而且我们看到该指令的确从常量内存中获取了其中一个操作数：
/*0050*/ IADD3 R0, R0, c[0x0][0x164], RZ ; /* 0x0000590000007a10 */
^^^^^^^^^^^^^
constant memory operand
因此负载之一（它是a或b，我不知道是哪一个，但是我猜是基于0x164的b）就在感兴趣的加法指令中。另一个负载也必须来自常量内存，那在哪里？我们还从上面的指令中注意到，“其他”操作数似乎已经存在于R0中（因为R0除了是目标操作数之外，它还是第一个源操作数）。因此，我们正在寻找一条指令，该指令从恒定内存（可能是“近”地址0x164）中提取内容，并将其放入R0中。在这里：
/*0020*/ IMAD.MOV.U32 R0, RZ, RZ, c[0x0][0x160] ; /* 0x00005800ff007624 */
因此，我们介绍了2个负载（均来自常量，因为它们是内核参数），添加和存储。
您可能会问，其余说明在做什么？我不会尝试解释每一个（因为我不能），但我认为它们分为2个一般类别：
创建所需的任何必要索引。该内核可能不需要很多索引，但是需要为R1创建一些索引，因为它用于创建本地存储操作的地址。这符合下面的项目2：
为printf函数调用设置“调用上下文”。此函数调用可能对寄存器，本地内存和堆栈中的数据有期望（GPU上的一种本地内存-R1可能有效地索引到堆栈中）。