找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 50|回复: 4

DAY28:

[复制链接]
发表于 2018-6-7 14:12:37 | 显示全部楼层 |阅读模式
5.2.3.1. Occupancy Calculator
Several API functions exist to assist programmers in choosing thread block size based on register and shared memory requirements.
  • The occupancy calculator API, cudaOccupancyMaxActiveBlocksPerMultiprocessor, can provide an occupancy prediction based on the block size and shared memory usage of a kernel. This function reports occupancy in terms of the number of concurrent thread blocks per multiprocessor.
    • Note that this value can be converted to other metrics. Multiplying by the number of warps per block yields the number of concurrent warps per multiprocessor; further dividing concurrent warps by max warps per multiprocessor gives the occupancy as a percentage.
  • The occupancy-based launch configurator APIs, cudaOccupancyMaxPotentialBlockSize and cudaOccupancyMaxPotentialBlockSizeVariableSMem, heuristically calculate an execution configuration that achieves the maximum multiprocessor-level occupancy.
The following code sample calculates the occupancy of MyKernel. It then reports the occupancy level with the ratio between concurrent warps versus maximum warps per multiprocessor.
  1. // Device code
  2. __global__ void MyKernel(int *d, int *a, int *b)
  3. {
  4.     int idx = threadIdx.x + blockIdx.x * blockDim.x;
  5.     d[idx] = a[idx] * b[idx];
  6. }

  7. // Host code
  8. int main()
  9. {
  10.     int numBlocks;        // Occupancy in terms of active blocks
  11.     int blockSize = 32;

  12.     // These variables are used to convert occupancy to warps
  13.     int device;
  14.     cudaDeviceProp prop;
  15.     int activeWarps;
  16.     int maxWarps;

  17.     cudaGetDevice(&device);
  18.     cudaGetDeviceProperties(&prop, device);
  19.    
  20.     cudaOccupancyMaxActiveBlocksPerMultiprocessor(
  21.         &numBlocks,
  22.         MyKernel,
  23.         blockSize,
  24.         0);

  25.     activeWarps = numBlocks * blockSize / prop.warpSize;
  26.     maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;

  27.     std::cout << "Occupancy: " << (double)activeWarps / maxWarps * 100 << "%" << std::endl;
  28.    
  29.     return 0;
  30. }
复制代码



The following code sample configures an occupancy-based kernel launch of MyKernel according to the user input.
  1. // Device code
  2. __global__ void MyKernel(int *array, int arrayCount)
  3. {
  4.     int idx = threadIdx.x + blockIdx.x * blockDim.x;
  5.     if (idx < arrayCount) {
  6.         array[idx] *= array[idx];
  7.     }
  8. }

  9. // Host code
  10. int launchMyKernel(int *array, int arrayCount)
  11. {
  12.     int blockSize;      // The launch configurator returned block size
  13.     int minGridSize;    // The minimum grid size needed to achieve the
  14.                         // maximum occupancy for a full device
  15.                         // launch
  16.     int gridSize;       // The actual grid size needed, based on input
  17.                         // size

  18.     cudaOccupancyMaxPotentialBlockSize(
  19.         &minGridSize,
  20.         &blockSize,
  21.         (void*)MyKernel,
  22.         0,
  23.         arrayCount);

  24.     // Round up according to array size
  25.     gridSize = (arrayCount + blockSize - 1) / blockSize;

  26.     MyKernel<<<gridSize, blockSize>>>(array, arrayCount);
  27.     cudaDeviceSynchronize();

  28.     // If interested, the occupancy can be calculated with
  29.     // cudaOccupancyMaxActiveBlocksPerMultiprocessor

  30.     return 0;
  31. }
复制代码


The CUDA Toolkit also provides a self-documenting, standalone occupancy calculator and launch configurator implementation in <CUDA_Toolkit_Path>/include/cuda_occupancy.h for any use cases that cannot depend on the CUDA software stack. A spreadsheet version of the occupancy calculator is also provided. The spreadsheet version is particularly useful as a learning tool that visualizes the impact of changes to the parameters that affect occupancy (block size, registers per thread, and shared memory per thread).














回复

使用道具 举报

 楼主| 发表于 2018-6-8 17:01:49 | 显示全部楼层
本章节主要说明了, 如何计算理论Occupancy。并给出了3种方式。

这三种方式各有特色。因为我们在之前的章节得知, occupancy往往和性能正相关(但不是100%绝对如此, 总有例外),所以往往对于任何一个kernel, 我们往往都需要尽量提高occupancy的。而提高occupancy就需要知道当前的occupancy是多少(一个0-1之间的数, 或者说0%到100%之间的数), 方能去有针对性的解决限制occupancy的因素。在实际的开发中, 我们会遭遇两种occupancy,一种叫理论occupancy,也就是能用手工(或者这里告诉你的各种API函数, 或者一个计算用的电子表格xlsx文件),另外一种是实际跑出来的,也就是在你的GPU上跑一遍kernel看看它能取得多少occupancy。后者叫achieved occupancy, 往往叫实际occupancy,或者取得的occupancy.
本章节说的这三种方式,均不是实际取得的occupancy。为了能有一个较好的对比,我将加上实际occupancy的方式在这里一起说,这样一共实际上4种方式。将会出现在本次阅读中。

首先说常规方式, 正常的方式我们往往总是在开发的机器上(例如一台GPU工作站,而不是GPU服务器),采用和目标部署时候(大规模运行的时候)相同的GPU。例如一个集群中,最终决定要部署的是Titan-V,那么往往工作站提供给开发用的也是Titan-V,虽然可能只有一两片,在这个机器上(或者说, 目标客户的机器上K80,我们开发的机器也用K80
)。这种开发方式是最常见的。因为最终结果如何(不考虑散热因素),可以在开发的时候基本上就知道了。此时在性能调优的时候,考虑到occupancy,是往往直接上profiiler,profiler会直接内置两种信息,一种是你当前的运行于卡的信息(也就是你目标部署时候的卡的信息),包括寄存器数量,Shared Memory数量,最大warps或者blocks数量限制,以及一些NV不想公开的因素。另外一种是profiler会直接知道你的Kernel的信息,包括kernel所使用的寄存器,shared memory,以及一些特殊的NV不想让你知道的信息。这样profiler会就地使用这两种信息,而完全不需要用户指定。此外,实际occupancy和运行时刻的状况有关,例如说,要被处理的数据导致性能上的变化,例如是否block中有特定的warps延迟或者提前结束。这些均会导致最终实际的occupancy被报告出来(很多时候基本等于理论occupancy,不等的时候往往只有一种。后文说)。而开发者立刻看到此信息后,可以选择调整自己的kernel,包括通过一些特殊的东西, launch bound, maxrreg数目,启动形状调整,shared memory大小调整,开发者立刻在此机器上进行上面这些调整(以及,一些特殊的调整),然后最终他会在开发的时刻就确定下来,一个最佳的资源使用,启动形状,和导致的occupancy。此时他会知道此时有最佳性能。

这种方式来确定occupancy和最佳性能的,是最常见的方式。但有的时候,这种方式不能用。为何?

  • 很多老板给员工配置的机器,和实际部署的机器的GPU不同。员工的开发的时候无法针对性的对此GPU进行调节优化。此时这种方式无法使用。
  • 在开发的时候,无法知道项目完成后的一定时间内, 例如3年,未来会应用到的新GPU会如何。此时无法直接在未来(现在还没有发布呢)的GPU上跑profiler的。
  • 现在是给广泛性的用户部署,用户可能有现在无法确定下来的配置情况。这种时候,开发时使用的profiler来现场确定的方式是无法生效的。



回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-6-8 17:59:55 | 显示全部楼层
此时需要此文中的另外3种方法。这3种方法各有特色。一种是使用excel电子表格。这个表格做的非常不错。
QQ图片20180608224900.png
如图,这个表格有3个步骤(左侧用户能改动的地方),例如我这里随意输入了:计算能力6.1(Pascal),只使用L2 cache,每个block需要128个线程每个线程需要48个寄存器,每个block需要4KB的shared memory,然后这个计算器会立刻告诉我(表格的内容会跟随更新)。
QQ图片20180608225148.png
当前我这种kernel的资源使用,将只能导致63%的occupancy,而且会告诉我(第二行红色)限制因素是因为我的寄存器使用太多了。然后我可以随意输入一点东西。
QQ图片20180608224934.png
如改成这种资源使用。则本表格内容会在下面自动改变:
QQ图片20180608224959.png
现在的限制因素变成了我使用了过小的block(只有32个线程的block),因为SM的block驻留数量有个固定的死限制。用32个这种只有1个warp的block会导致只有50%的occupancy,这点其实之前说过。现在常见的主流卡至少需要使用64的block大小,也就是一个block里至少需要有2个warps,才有可能压满设备。
实验不同的数据,例如我刚才说的block大小数据,然后你可以立刻在改成32的时候,知道这个信息。这个计算器是本章节最后说的那个东西,路径也在本章节里的英文里。默认在你的C盘的CUDA开发包安装目录的某个子目录(子文件夹)下。这个电子表格是本章节的3种方式的最后一种。
你可以看到它的优势:1)知道了自己的kernel的使用资源情况后,可以手工填入。不需要你有计算能力X.Y的卡,不需要这卡在手边,就可以直接在这个表格里设定,而表格也会直接告诉你该卡的下的理论情况。非常方便。这样适合那种开发的时候卡是真没有的人。弥补了基于profiler的开发流程中必须需要有实体该卡的缺陷。此外,如果之前说的,该表格还告诉了你一些手册从来不说的东西,例如寄存器的分配粒度,Shared Memory的分配粒度。例如很多人以为我的kernel原本使用19个寄存器,现在我改成使用17个寄存器,然后这个表格后续单独的一个“GPU Data"的表单会告诉你,并没有。因为某些计算能力的卡上面,寄存器的分配粒度是8个,只有8个,16个,24个,32个。。。这种,你使用了17个和19个,实际都会使用24个的。这种信息就规避了盲目的去优化一些东西。



QQ图片20180608224822.png
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-6-8 21:41:00 | 显示全部楼层
因为没有必要。优化了也没有结果。而这些是手册不说的,本表格是直接的一手信息来源(NV),使用该表格需要知道一些信息,例如kernel使用了多少寄存器,这个信息可以在编译的时候通过ptxas的详细输出得到,手册的后面会提到如何得到这个信息。这也是和之前的用profiler获取实际occupancy的不同,该表格只会自动知道设备的信息(例如计算能力6.1的1080),而不能知道你的kernel函数的信息,因为它只是一个普通的表格(好无辜的感觉),但依然已经比你在没有任何卡的情况下,纯靠蒙的方式要好。

本章节还提供了另外一种方式,叫做通过CUDA提供的API。这种方式比表格更加方便。主要有2个。(1)可以让程序知道自己的信息。程序通过1个或者多个API调用,可以在运行的时刻(例如在用户的机器上,显卡是K80的时候)知道目前的卡上面情况。而你却不能让用户随身带着一本电子表格。自己去调。

程序知道了自己在这个卡上的情况后,可以自己实验不同的occupancy组合(例如在启动热身阶段,或者一次性的在安装程序阶段),从而达到程序中目标用户机器上的性能最优化。(2)只需要少量的信息。比使用电子表格少。因为一个CUDA项目在实际的运行的时候,一些信息可以自动从自身获得,例如卡的信息,例如kernel的固定资源使用(寄存器,固定的shared memory),此时只需要提供少量的例如启动block形状,动态shared memory之类的少量信息,即可获取到在可能用这个少量的动态信息启动kernel后,kernel会有什么样子的预估occupancy(然后程序可以决定自己是否真的用这种形状启动。或者尝试其他形状组合之类),这样的话,程序可以在运行的时候自我调优(这个实际上是个交复杂的过程,因为occupancy不等于一切,但本章节的方式很多时候是够用的),这种方式只需要用户提供一些少量的动态信息即可,得到的也是理论结果。比电子表格需要输入的手工大量信息要少。但比profiler的全自动要多,处于折中。

你可以看一下本章节的具体代码的使用,比如cudaOccupancy开头的那些函数,如何从blocks数量变换到warps数量(乘以block里的线程数除以32即可),如何从warps数量变换到occupancy等等。
Note that this value can be converted to other metrics,

这就是这里说的。

Multiplying by the number of warps per block yields the number of concurrent warps per multiprocessor; further dividing concurrent warps by max warps per multiprocessor gives the occupancy as a percentage.


这些。
yield这里是得到/产出的意思。
为何这种方式能使用较少的信息,是因为CUDA运行的时候自己能得到一些。这里是显然的。
然后还有另外一个最后的方式,类似这种,也是在运行的时候能自我计算occupancy,叫做cuda occ(注意不是cuda occupancy),也就是一个叫cuda_occupancy.h的头文件(这文件里有一些被inline的代码),使用该头文件有个最大的好处。就是不需要一张支持CUDA的卡(刚才那种需要的),例如我可以在A卡上跑代码,然后同时给用户比照一下在N卡上会如何。
该头文件的最大好处是没有依赖项。方便集成在任何项目中。
例如一个集群中可能会有一些负责任务调度的机器。该机器可能具体的会发布任务的时候细节到具体的提供kernel启动信息。而这个机器因为自己不需要使用CUDA的,它完全可能可以没有卡。此时该头文件(cuda occ)依然可以运行。此外,该头文件的最大好处是,里面提供了代码,会告诉你为何一些东西,影响occupancy,又是如何影响的。看了这个代码比使用profiler,使用电子表格,使用cudaOccupancy*(),都有好处。前3个只知道结果,不知所以然。而使用这个文件会知道为何(打开看一下就可以了),适合需要进阶的用户,或者手头完全无卡的用户。

回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-6-8 22:02:39 | 显示全部楼层
关于profiler还有很多要说的。但不妨等后文遇到再说。但这里必须要提到的是:计算能力5.2+(包括计算能力5.3, TX1躺枪)和Pascal 6.1,使用L1 cache会导致有的时候无法启动kernel(真的),具体原因尚未明确(NV没说。这里我们也不说),当在这些设备上使用要求使用L1 cache(Unified cache)的kernel的时候,如果对L1 Cache的使用会导致kernel启动失败,(occupancy为0),那么CUDA会自动以禁用L1的方式启动你的kernel。其实会导致意想不到的性能影响(不一定是正面还是负面),cudaOccupancy*()系列函数(第三种方式),会自动默认当可能导致occupancy为0%的时候自动为你修补(以禁用L1的方式为你报告结果),但cudaOccupancy*()有个特定的参数(看Runtime API手册),会告诉他不要自动修补。此时指定此参数会真实的报告0%的occupancy。
此时你会知道此时无法使用L1。目前尚未知道此特性在这些卡上(还不清楚6.2如何。不能知道TX2如何。TX1是有这个问题的),这是一个特性,还是一个BUG?(看上去更像是一个特性)

手册只提到了(就在这个手册后面)GM204有此问题(GTX980, GTX970。或者对应的Quadro),但实际上所有的计算能力5.2/5.3/6.1的卡均有此问题。NV目前没有在手册里具体说明此问题。


只有一句话:
QQ图片20180608224822.png

在GM204上启用(一级)缓存会影响occupancy。如果某block在启用缓存和一定资源的时候会导致0%的occupancy,那么CUDA Driver会自动改掉此缓存设定,以便运行kernel能成功启动。profilr会报告此情况。然后无更多解释。第三方来源有更详细的解释信息, 但这里就不提供了。(注意,不仅仅是GM204。此问题在Pascal(6.1)上依然存在)用户需要注意这点。有趣的是,GM107(第一代maxwell)上并无此问题。







回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表