为什么 CUDA 内核的开头有一个未使用的数据移动？

Question

我正在尝试研究SASS从非常基本的 CUDA 内核生成的文件。这是内核：

__global__ void kernel(const float * x,
                       float * y,
                       const uint num_rows,
                       const uint num_cols) {
    const uint num_elems = num_rows * num_cols;
    const uint tid = blockDim.x * blockIdx.x + threadIdx.x;
    for (uint idx = tid; idx < num_elems; idx += blockDim.x * gridDim.x) {
        y[idx] = x[idx];
    }
}

这是SASS文件。

1   00007f26 14f69f00         MOV R1, c[0x0][0x28]
2   00007f26 14f69f10         S2R R0, SR_CTAID.X
3   00007f26 14f69f20         ULDC.64 UR4, c[0x0][0x178]
4   00007f26 14f69f30         UIMAD UR4, UR5, UR4, URZ
5   00007f26 14f69f40         S2R R3, SR_TID.X  3   3840
6   00007f26 14f69f50         IMAD R0, R0, c[0x0][0x0], R3
7   00007f26 14f69f60         ISETP.GE.U32.AND P0, PT, R0, UR4, PT
8   00007f26 14f69f70   @P0   EXIT
9   00007f26 14f69f80         ULDC.64 UR6, c[0x0][0x118]
10  00007f26 14f69f90         MOV R5, 0x4
11  00007f26 14f69fa0         IMAD.WIDE.U32 R2, R0, R5, c[0x0][0x160]
12  00007f26 14f69fb0         LDG.E R3, [R2.64]
13  00007f26 14f69fc0         IMAD.WIDE.U32 R4, R0, R5, c[0x0][0x168]
14  00007f26 14f69fd0         MOV R7, c[0x0][0x0]                               
15  00007f26 14f69fe0         IMAD R0, R7, c[0x0][0xc], R0
16  00007f26 14f69ff0         ISETP.GE.U32.AND P0, PT, R0, UR4, PT
17  00007f26 14f6a000         STG.E [R4.64], R3
18  00007f26 14f6a010   @!P0  BRA 0x7f2614f69f90
19  00007f26 14f6a020         EXIT
20  00007f26 14f6a030         BRA 0x7f2614f6a030                            

问题：

在第一行中SASS，c[0x0][0x28]被转移到R1，并且我们从未使用过它。此行为不限于此内核。我已经用几个不同的简单内核对其进行了测试，并且总是看到这条指令。有谁知道这条指令的目的是什么？

更多信息：

1. 我们知道 c[0x0][xyzw]（因此，bank 0x0）存储到内核参数和启动配置。不过，目前还不清楚为什么第一行会有一个看似无用的举动。
2. 这不是一个大文件的一小部分。你看到的就是我编译的。

Answer 1

我也没有找到这方面的文档。然而，R1似乎充当堆栈指针。您可以在如下代码中看到它的使用：

__global__ void foo(int* inout) {
    int tid = threadIdx.x;
    volatile int local[12];
    local[inout[tid]] = 12;
    inout[tid] = local[inout[tid + 1]];
}

foo(int*):
 MOV R1, c[0x0][0x20] 
 IADD32I R1, R1, -0x30 
 S2R R4, SR_TID.X         
 SHR R0, R4.reuse, 0x1e 
 ISCADD R4.CC, R4, c[0x0][0x140], 0x2 
 IADD.X R5, R0, c[0x0][0x144] 
 LDG.E R0, [R4] 
 LDG.E R2, [R4+0x4] 
 MOV32I R3, 0xc 
 LEA R0, R0, R1.reuse, 0x2 
 LEA R2, R2, R1, 0x2 
 STL [R0], R3         
 LDL R2, [R2] 
 STG.E [R4], R2 
 NOP 
 EXIT 
.L_x_0:
 BRA `(.L_x_0) 
 NOP
.L_x_1:

常量内存中初始偏移量的位置似乎在不同架构之间发生变化，因为这里（SM 5.2）它是 0x20。

但是，它似乎不是字面量堆栈指针，因为所有线程都以相同的值开始。和指令会考虑一些每个线程的偏移量和比例，以便只要 warp 中的所有线程访问相同的相对地址，32 位访问就会完全STL合并。LDL

至于为什么这个负载没有消除，我不知道。也许调试器或其他一些机制（例如机器异常处理）始终具有有效的堆栈指针。