6.7 并发内核处理

SM 2.x架构和更高级别的GPU能够并发运行多个内核，只要它们是在不同流中被启动的并且有适当大小的内存块（足够小，从而一个内核不会填满整个GPU）的话。假使每个内核启动中线程块的数量足够小的话，代码清单6-5（9～14行）将导致内核并发运行。由于内核之间只能通过全局内存通信，我们可以添加以下代码到AddKernel()来跟踪同时运行的内核的数量。可以使用下面的“内核并发跟踪”结构。

static const int g_maxStreams = 8;  
typedef struct KernelConcurrencyData_st{int mask; // mask of active kernelsint maskMax; // atomic max of mask popcountint masks[g_maxStreams];int count; // number of active kernelsint countMax; // atomic max of kernel countint counts[g_maxStreams];} KernelConcurrencyData;

我们可以在函数的开始和结尾分别添加代码到AddKernel()中来“登记”和“退房”。“登记”使用“内核ID”参数kid（一个传递给内核的值，其范围介于0到NumStreams-1之间），接着计算一个对应一个全局变量的内核ID的掩膜（kid左移一位），并对该值进行原子或操作存入全局变量里。请注意，atomicOR()在OR执行之前返回在内存该位置的值。结果，当原子或操作执行时，返回值对每一个活跃的内核都有一个位来对应。

同样地，这段代码通过递增 Counts->count 和调用在共享的全局内存上的 atomicMax() 来追踪活跃内核的数量。

// check in, and record active kernel mask and count  
// as seen by this kernel.  
if (_kernelData && blockIdx.x == 0 && threadIdx.x == 0) {  
    int myMask = atomicOr( &kernelData->mask, 1 << kid);  
    kernelData->masks[kid] = myMask | (1 << kid);  
    int myCount = atomicAdd( &kernelData->count, 1 );  
    atomicMax( &kernelData->countMax, myCount + 1 );  
    kernelData->counts[kid] = myCount + 1;  
}

在内核的底部，类似的代码清除掩膜并递减活跃内核的计数。

// check out  
if (tokenizerData && blockIdx.x==0 && threadIdx.x==0) {  
    atomicAnd( &tokenizerData->mask, -(1<<kid));  
    atomicAdd( &tokenizerData->count, -1);  
}

kernelData参数代表在文件作用域内声明的__device__变量。

device_KernelConcurrencyData g_kernelData;

请记住，指向g_kernelData的指针必须通过调用CUDAGetSymbolAddress()获得。可以编写引用&g_kernelData的代码，但是CUDA的语言集成特性将无法正确解析该地址。

concurrencyKernelKernel.cu程序增加了对命令行选项blocksPerSM的支持，来指定内核启动的线程块的数量。它将生成活跃内核数量的报告。两个concurrencyKernelKerne的调用示例如下。

$ ./concurrencyKernelKernel -blocksPerSM 2
Using 2 blocks per SM on GPU with 14 SMs = 28 blocks
Timing sequential operations... Kernel data:
Masks: (0x1 0x0 0x0 0x0 0x0 0x0)
Up to 1 kernels were active: (0x1 0x0 0x0 0x0 0x0 0x0)
Timing concurrent operations...
Kernel data:
Masks: (0x1 0x3 0x7 0xe 0x1c 0x38 0x60 0xe0)
Up to 3 kernels were active: (0x1 0x2 0x3 0x3 0x3 0x2 0x3)
$ ./concurrencyKernelKernel -blocksPerSM 3
Using 3 blocks per SM on GPU with 14 SMs = 42 blocks
Timing sequential operations... Kernel data:
Masks: (0x1 0x0 0x0 0x0 0x0 0x0 0x0)
Up to 1 kernels were active: (0x1 0x0 0x0 0x0 0x0 0x0 0x0)
Timing concurrent operations... Kernel data:
Masks: (0x1 0x3 0x6 0xc 0x10 0x30 0x60 0x80)
Up to 2 kernels were active: (0x1 0x2 0x2 0x2 0x1 0x2 0x1)

需要注意的是，blocksPerSM是每个内核启动的块数，所以numStreams*blocksPerSM个块是在numStreams个独立内核里启动的。你可以看到，当内核网格较小时，硬件可以并发运行更多的内核。但是对于本章讨论的工作负载，并发内核处理没有带来性能上的好处。

6.8 GPU/GPU同步：CUDAWaitEvent()

到现在为止，本章描述的所有同步函数都已经和CPU/GPU同步相关联。它们或者等待或者查询GPU操作的状态。CUDAStreamWaitEvent()函数是与CPU异步的，使指定的stream等待直到事件被记录。流和事件不必与同一个CUDA设备相关联。9-3节介绍了如何进行GPU之间的同步，并且使用了这一特性来实现一个点对点的内存复制（见代码清单9-1）。

多GPU的流和事件：注意事项和限制

·流和事件存在于上下文（或设备）的作用域内。当调用

cuCtxDestroy()或CUDADevice Reset()时，相关联的流和事件都会被破坏。

·内核启动和cu(da)EventRecord()仅仅能在同一个上下文/设备中

使用CUDA流。

· CUDAMemcpy()可以从任何流调用，但最好是源上下文/设备中调

用。

· CUDAWaitEvent()可以使用任何流、被任何事件调用。