网格同步(Grid Synchronization):
在引入协作组之前,CUDA编程模型只允许在内核完成边界上的线程块之间进行同步。 内核边界带有隐含的状态无效,并带有潜在的性能影响。
例如,在某些使用情况下,应用程序拥有大量的小内核,每个内核代表处理管道中的一个阶段。 当前CUDA编程模型需要这些内核的存在,以确保在一个流水线阶段上运行的线程块在下一个流水线阶段运行的线程块准备好使用之前生成数据。 在这种情况下,提供全局线程间块同步的能力将允许应用程序重新构建为具有持久线程块,当给定阶段完成时,这些线程块能够在设备上进行同步。
要在内核中同步整个网格,您只需使用该组:
grid_group grid = this_grid();
调用:
grid.sync();
要启用网格同步,启动内核时,需要使用cuLaunchCooperativeKernel CUDA运行时启动API,而不是<<< … >>>执行配置语法:
cudaLaunchCooperativeKernel(
const T *func,
dim3 gridDim,
dim3 blockDim,
void **args,
size_t sharedMem = 0,
cudaStream_t stream = 0
)
(或等效的CUDA驱动程序)。
为了保证GPU上的线程块的共同驻留,需要仔细考虑启动的块的数量。 例如,每个SM的块可以按如下方式启动:
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
// initialize, then launch
cudaLaunchCooperativeKernel((void*)my_kernel, deviceProp.multiProcessorCount,
numThreads, args);
或者,您可以使用占用率计算器计算每SM可同时容纳多少个块,如下所示:
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm, my_kernel,
numThreads, 0));
// initialize, then launch
cudaLaunchCooperativeKernel((void*)my_kernel, numBlocksPerSm, numThreads, args);
还要注意,要使用网格同步,必须在单独的编译中编译器件代码(请参阅CUDA编译器驱动程序NVCC文档中的“在CUDA中使用独立编译”一节)以及链接的器件运行时。最简单的示例是:
nvcc -arch=sm_61 -rdc=true mytestfile.cu -o mytest
您还应该确保设备支持协作启动属性,这可以通过使用cuDeviceAttribute CUDA驱动程序API来确定:
int pi=0;
cuDevice dev;
cuDeviceGet(&dev,0) // get handle to device 0
cuDeviceGetAttribute(&pi, CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH, dev);
如果在设备0上支持该属性,将会将pi设置为1。
多设备同步:
为了启用多个设备与协同组之间的同步,需要使用cuLaunchCooperativeKernelMultiDevice CUDA API。 这与现有的CUDA API显着不同,将允许单个主机线程跨多个设备启动内核。 除了cuLaunchCooperativeKernel所做的约束和保证外,这个API还有其他的语义:
__device__
,__constant__
或__managed__
设备全局变量将在每个设备上独立实例化。 用户负责适当地初始化这些设备全局变量。启动参数应该使用struct定义:
typedef struct CUDA_LAUNCH_PARAMS_st {
CUfunction function;
unsigned int gridDimX;
unsigned int gridDimY;
unsigned int gridDimZ;
unsigned int blockDimX;
unsigned int blockDimY;
unsigned int blockDimZ;
unsigned int sharedMemBytes;
CUstream hStream;
void **kernelParams;
} CUDA_LAUNCH_PARAMS;
并传入启动API:
cudaLaunchCooperativeKernelMultiDevice(
CUDA_LAUNCH_PARAMS *launchParamsList,
unsigned int numDevices);
以类似于上述网格范围同步的方式。 另外,与网格同步一样,生成的设备代码看起来非常相似:
multi_grid_group multi_grid = this_multi_grid();
multi_grid.sync();
并需要在单独的编译中进行编译。
您还应该确保设备支持协作式多设备启动属性,其方式与上一节中所述类似,但使用CU_DEVICE_ATTRIBUTE_COOPERATIVE_MULTI_DEVICE_LAUNCH而不是CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH。
手机扫一扫
移动阅读更方便
你可能感兴趣的文章