CUDA内核执行与流式多处理器详解

### CUDA内核执行与流式多处理器详解 #### 1. 线程ID与warp映射 在14Wx8H的块中，每个warp仅容纳28个线程，在整个内核执行过程中，有12.5%的线程通道处于空闲状态。为避免这种资源浪费，开发者应确保块包含32的倍数个线程。 #### 2. 占用率（Occupancy） 占用率是衡量在给定内核启动时，每个流式多处理器（SM）上实际运行的线程数与该SM上潜在可运行的最大线程数之比。计算公式为： \[占用率 = \frac{每个SM的warp数}{每个SM的最大warp数}\] 分母（每个SM的最大warp数）仅取决于设备的计算能力，是一个常数。而分子（每个SM的warp数），即决定占用率的因素，受以下参数影响： - 计算能力（1.0, 1.1, 1.2, 1.3, 2.0, 2.1, 3.0, 3.5） - 每个块的线程数 - 每个线程的寄存器数 - 共享内存配置 - 每个块的共享内存 为帮助开发者评估这些参数之间的权衡，CUDA工具包提供了一个以Excel电子表格形式存在的占用率计算器。输入上述参数后，该电子表格将计算以下结果： - 活跃线程数 - 活跃warp数 - 活跃块数 - 占用率（活跃warp数除以硬件的最大活跃warp数） 同时，电子表格还会指出限制占用率的参数，包括： - 每个多处理器的寄存器数 - 每个多处理器的最大warp或块数 - 每个多处理器的共享内存 需要注意的是，占用率并非CUDA性能的唯一决定因素。有时，每个线程使用更多的寄存器并依靠指令级并行性（ILP）来提升性能可能是更好的选择。 例如，对于某些低占用率的内核，也能实现接近最大的全局内存带宽。如GlobalReads内核的内循环可以根据模板参数进行展开，随着展开迭代次数的增加，所需的寄存器数量增加，占用率下降。以cg1.4xlarge实例类型中的Tesla M2050为例，在禁用ECC的情况下，报告的峰值读取带宽为124GiB/s，占用率为66%。 #### 3. 动态并行性（Dynamic Parallelism） 动态并行性是一项仅适用于SM 3.5级硬件的新功能，它允许CUDA内核启动其他CUDA内核，并调用CUDA运行时的各种函数。使用动态并行性时，CUDA运行时的一个子集（即设备运行时）可供设备上运行的线程使用。 动态并行性引入了“父”和“子”网格的概念。由另一个CUDA内核调用的内核（与之前所有CUDA版本中由主机代码调用不同）是“子内核”，调用它的网格是其“父”。默认情况下，CUDA支持两级嵌套（一级为父级，一级为子级），可以通过调用`cudaSetDeviceLimit()`并传入`cudaLimitDevRuntimeSyncDepth`来增加嵌套级别。 动态并行性旨在解决那些之前需要将结果返回给CPU，以便CPU指定在GPU上执行哪些工作的应用程序。这种“握手”过程会破坏CPU/GPU执行管道中的并发性能，而动态并行性通过使GPU能够从内核中自行启动工作，避免了这些管道气泡。 动态并行性在以下几种情况下可以提高性能： - 能够在核函数开始执行之前初始化所需的数据结构，而之前这种初始化必须在主机代码中完成或通过调用单独的内核来处理。 - 为诸如Barnes - Hut引力积分或空气动力学模拟的分层网格评估等应用程序实现简化的递归。 需要注意的是，动态并行性仅在单个GPU内有效，内核可以调用内存复制或其他内核，但不能将工作提交到其他GPU。 #### 3.1 作用域和同步（SCOPING AND SYNCHRONIZATION） 除了块和网格大小外，子网格继承了其父网格的大多数内核配置参数，如共享内存配置（由`cudaDeviceSetCacheConfig()`设置）。线程块是作用域的单位，线程块创建的流和事件只能由该线程块使用，并且在线程块退出时会自动销毁。 CUDA保证父网格在其所有子网格完成之前不会被视为完成。尽管父网格可以与子网格并发执行，但直到父网格调用`cudaDeviceSynchronize()`，子网格才会开始执行。 如果一个线程块中的所有线程都退出，该线程块的执行将暂停，直到所有子网格完成。如果这种同步不够，开发者可以使用CUDA流和事件进行显式同步。在设备运行时，流和事件的使用有一些限制： - 流和事件只能在创建它们的线程块内使用。 - NULL流在设备运行时的语义与主机运行时不同，在设备上，NULL流是独立的流，任何流间同步都必须使用事件来执行。 - 创建流时必须传递`cudaStreamNonBlocking`标志，`cudaStreamSynchronize()`调用不受支持，同步必须通过事件和`cudaStreamWaitEvent()`来实现。 - 仅支持CUDA事件的流间同步功能，`cudaEventSynchronize()`、`cudaEventElapsedTime()`和`cudaEventQuery()`不受支持，并且由于不支持计时，事件必须通过传递`cudaEventDisableTiming`标志来创建。 #### 3.2 内存模型（MEMORY MODEL） 父网格和子网格共享相同的全局和常量内存存储，但它们有各自独立的本地和共享内存。 ##### 全局内存（Global Memory） 在子网格的执行过程中，有两个时间点其内存视图与父网格完全一致：当子网格被父网格调用时，以及当子网格完成（由父线程中的同步API调用发出信号）时。 父线程在调用子线程之前的所有全局内存操作对子网格可见。在父网格同步子网格完成后，子网格的所有内存操作对父网格可见。零拷贝内存与全局内存具有相同的一致性和连贯性保证。 ##### 常量内存（Constant Memory） 常量是不可变的，在核函数执行期间不能从设备进行修改。在核函数线程中获取常量内存对象的地址与所有CUDA程序具有相同的语义，并且在父网格和子网格之间传递该指针是完全支持的。 ##### 共享和本地内存（Shared and Local Memory） 共享内存和本地内存分别对线程块或线程是私有的，在父网格和子网格之间不可见或不连贯。当在其作用域之外引用这些位置的对象时，行为是未定义的，可能会导致错误。 如果`nvcc`检测到试图误用指向共享或本地内存的指针，它将发出警告。开发者可以使用`__isGlobal()`内建函数来确定给定指针是否引用全局内存。指向共享或本地内存的指针不是`cudaMemcpy*Async()`或`cudaMemset*Async()`的有效参数。 本地内存是执行线程的私有存储，在该线程之外不可见。在启动子内核时，传递指向本地内存的指针作为启动参数是非法的，子内核解引用这样的本地内存指针的结果是未定义的。为确保编译器不会无意中违反