CUDA编程细节大杂烩

编程相关
1

细节1

for (unsigned int i = 0; i < MEMCOPY_ITERATIONS; i++) {
  checkCudaErrors(
    cudaMemcpy(d_odata, d_idata, memSize, cudaMemcpyDeviceToDevice));
}

bandwidthInGBs = (2.0f * memSize * (float)MEMCOPY_ITERATIONS) / (double)1e9，在官方示例中这么计算带宽，乘以2是因为考虑了数据的读取和写入两个步骤。在 GPU 内存传输中，即使是单向传输，也涉及到两个过程：将数据从源内存位置读取（d_idata ），然后将数据写入目标内存位置（d_odata ）。

细节2

一个较为微妙的观点是，每次访问寄存器所涉及的指令数量比访问全局内存要少。在大多数现代处理器中，算术指令具有“内置”的寄存器操作数。例如，浮点加法指令可能是以下形式：

fadd r1, r2, r3

其中 r2 和 r3 是寄存器编号，指定了寄存器文件中可以找到输入操作数值的位置。存储浮点加法结果值的位置由 r1 指定。因此，当算术指令的一个操作数在寄存器中时，不需要额外的指令来使操作数值可用于算术和逻辑单元（ALU），在那里进行算术计算。

与此同时，如果一个操作数值在全局内存中，处理器需要执行一个内存加载操作，以使操作数值可用于 ALU。例如，如果浮点加法指令的第一个操作数在全局内存中，所涉及的指令可能看起来如下示例：

load r2, r4, offset
fadd r1, r2, r3

其中加载指令将一个偏移值加到 r4 的内容上，形成操作数值的地址。然后它访问全局内存，并将该值放入寄存器 r2 中。一旦操作数值在 r2 中，fadd 指令使用 r2 和 r3 中的值执行浮点加法，并将结果放入 r1 中。由于处理器每个时钟周期只能获取和执行有限数量的指令，所以包含额外加载的版本可能需要更多时间来处理。这是将操作数放入寄存器可以提高执行速度的另一个原因。

细节3

在cuda kernel访问tensor的时候要记得确保tensor是连续的，比如使用torch_check contiguous，（tensor slice之后传进了kernel，这种bug不会报错，但就是计算结果不对）

速度特别快的kernel

在CUDA流中的不同内核似乎无法并行运行。原帖作者通过一个示例程序，试图在不同的CUDA流中并行运行几个内核，但观察到这些内核并没有如预期那样并行执行。回帖中，有人提出可能是因为这些内核的执行速度过快，以至于在一个内核结束之前，下一个内核还没有开始运行的机会。另一个回帖指出，这些内核由于没有副作用，被编译器优化为了空内核（null kernel）。原帖作者通过修改内核代码，增加了计算量，并避免了编译器优化，最终实现了预期的并行执行效果。

https://forums.developer.nvidia.com/t/kernels-in-cuda-streams-seems-not-running-in-parallel/286549

参考