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。
timg

目录
相关文章
|
并行计算 C语言 编译器
|
移动开发 并行计算 编译器
|
并行计算 API 调度
CUDA学习(八十八)
3.虽然__syncthreads()一直被记录为同步线程块中的所有线程,但Pascal和以前的体系结构只能在warp级别强制执行同步。 在某些情况下,只要每条经线中至少有一条线达到屏障,就可以在不被每条线执行的情况下成功实现屏障。
1691 0
|
并行计算 前端开发 Windows
|
并行计算 编译器
|
并行计算 C语言 编译器
|
并行计算 API C语言
|
并行计算 编译器 存储
|
并行计算 API 编译器
CUDA学习(六十五)
很早之前就发现云栖社区的编辑器有一个Bug,往草稿箱存博客,当草稿箱博客数超过十篇时,无法再选择十篇前的博客进行编辑
2364 0
|
并行计算 API 异构计算