CUDA实践指南（二十二）-白红宇

并发内核执行：

正如“计算异步和重叠传输”中所述，CUDA流可用于将内核执行与数据传输重叠。在能够并发执行内核的设备上，还可以使用流同时执行多个内核，以更充分地利用设备的多处理器。设备是否具有此功能由cudaDeviceProp结构的并发内核字段指示（或在deviceQuery CUDA示例的输出中列出）。非默认流（流0以外的流）对于并发执行是必需的，因为使用默认流的内核调用只有在设备上的所有前面的调用（任何流）完成之后才开始，并且设备上没有任何操作流）开始直到完成。

以下示例说明了基本技巧。因为kernel1和kernel2在不同的非默认流中执行，所以有能力的设备可以同时执行内核。

cudaStreamCreate(&stream1);cudaStreamCreate(&stream2);kernel1<<
    
     >>(data_1);kernel2<<
     
      >>(data_2);

多个上下文：

CUDA工作发生在称为上下文的特定GPU的进程空间内。上下文封装了该GPU的内核启动和内存分配以及支持页表等内容。上下文在CUDA Driver API中是明确的，但在CUDA Runtime API中完全隐含，它自动创建和管理上下文。

借助CUDA驱动程序API，CUDA应用程序进程可能会为给定的GPU创建多个上下文。如果多个CUDA应用程序进程同时访问同一GPU，则这几乎总是暗含多个上下文，因为上下文与特定主机进程绑定，除非正在使用CUDA多进程服务。

虽然可以在给定的GPU上同时分配多个上下文（及其相关资源，如全局内存分配），但只有其中一个上下文可以在该GPU上的任何给定时刻执行工作; 共享相同GPU的上下文是时间片的。创建额外的上下文会导致每上下文数据的内存开销和上下文切换的时间开销。此外，当需要并行执行多个上下文时，对上下文切换的需求可以降低利用率。

因此，最好在相同的CUDA应用程序中避免每个GPU的多个上下文。为此，CUDA Driver API提供了访问和管理每个GPU上称为主要上下文的特殊上下文的方法。这些是CUDA运行时隐式使用的相同上下文，当没有线程的当前上下文时。

// When initializing the program/libraryCUcontext ctx;cuDevicePrimaryCtxRetain(&ctx, dev);// When the program/library launches workcuCtxPushCurrent(ctx);kernel<<<...>>>(...);cuCtxPopCurrent(&ctx);// When the program/library is finished with the contextcuDevicePrimaryCtxRelease(dev);

NVIDIA®（英伟达™）SMI可用于配置独立处理模式的GPU，从而将每个GPU的上下文数量限制为一个。在创建过程中，此上下文可以是任意多个线程，如果在设备上已经存在使用CUDA驱动程序API创建的非主要上下文，cuDevicePrimaryCtxRetain将会失败。

隐藏寄存器依赖关系:

中等优先级：为了隐藏由寄存器依赖性引起的延迟，为每个多处理器维护足够数量的活动线程（即足够的占用率）。

当指令使用存储在指令写入的寄存器中的结果之前，会产生寄存器相关性。当前支持CUDA的GPU的延迟时间大约为24个周期，因此线程在使用算术结果之前必须等待24个周期。但是，这种延迟可以通过执行其他warp中的线程完全隐藏。