CUDA编程细节大杂烩

细节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不会报错,但就是计算结果不对)