CUDA学习(五十七)
网格同步(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还有其他的语义:
- 该API将确保启动是原子级的,即,如果API调用成功,则提供的线程块数将在所有指定的设备上启动。
- 通过此API启动的功能必须相同。 驱动在这方面没有进行明确的检查,因为这在很大程度上是不可行的。 确保这一点取决于应用程序。
- 提供的launchParamsList中没有两个条目可映射到同一设备。
- 本次发布所针对的所有设备必须完全相同。 即它们必须具有相同的主号码和次号码。
- 所有设备的块大小,网格大小和每个网格的共享内存数量必须相同。 请注意,这意味着每个设备可以启动的块的最大数量将受SM数量最少的设备的限制。
- 拥有正在启动的CU函数的模块中存在的任何用户定义的
__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。
推荐阅读
-
CUDA学习(五十七)
-
kubernetes(k8s) 学习 (十) k8s存储之 Volumes (简介+emptyDir卷+hostpath卷+nfs卷)
-
Linux学习之CentOS7安装Mysql5.7直接覆盖Mariadb
-
有关于node缓冲区的学习
-
HTML&CSS基础学习笔记1.13—无序列表 博客分类: HTML
-
HTML&CSS基础学习笔记1.28-给网页添加一个css样式 博客分类: HTML
-
C语言基础学习学习前的准备-1 博客分类: C语言 C语言位字字节编程习惯
-
Python新手学习基础之循环结构——For语句 博客分类: Python python循环语句for遍历
-
Java 学习笔记- I/O - Byte Stream 博客分类: Java 学习笔记 javaiObyte stream字节流
-
javaIO学习之字节流和字符流 博客分类: j2se javaio字节流字符流