DAY28:

  • 4 replies
  • 69 views
*

sisiy

  • *****
  • 121
    • 查看个人资料
DAY28:
« 于: 六月 07, 2018, 02:12:37 pm »
5.2.3.1. Occupancy CalculatorSeveral 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.[indent]
程序代码: [选择]
// Device code
__global__ void MyKernel(int *d, int *a, int *b)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    d[idx] = a[idx] * b[idx];
}

// Host code
int main()
{
    int numBlocks;        // Occupancy in terms of active blocks
    int blockSize = 32;

    // These variables are used to convert occupancy to warps
    int device;
    cudaDeviceProp prop;
    int activeWarps;
    int maxWarps;

    cudaGetDevice(&device);
    cudaGetDeviceProperties(&prop, device);
   
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &numBlocks,
        MyKernel,
        blockSize,
        0);

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

    std::cout << "Occupancy: " << (double)activeWarps / maxWarps * 100 << "%" << std::endl;
   
    return 0;
}



The following code sample configures an occupancy-based kernel launch of MyKernel according to the user input.
程序代码: [选择]
// Device code
__global__ void MyKernel(int *array, int arrayCount)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < arrayCount) {
        array[idx] *= array[idx];
    }
}

// Host code
int launchMyKernel(int *array, int arrayCount)
{
    int blockSize;      // The launch configurator returned block size
    int minGridSize;    // The minimum grid size needed to achieve the
                        // maximum occupancy for a full device
                        // launch
    int gridSize;       // The actual grid size needed, based on input
                        // size

    cudaOccupancyMaxPotentialBlockSize(
        &minGridSize,
        &blockSize,
        (void*)MyKernel,
        0,
        arrayCount);

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

    MyKernel<<>>(array, arrayCount);
    cudaDeviceSynchronize();

    // If interested, the occupancy can be calculated with
    // cudaOccupancyMaxActiveBlocksPerMultiprocessor

    return 0;
}


The CUDA Toolkit also provides a self-documenting, standalone occupancy calculator and launch configurator implementation in /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).













[/indent]

*

sisiy

  • *****
  • 121
    • 查看个人资料
(无标题)
« 回复 #1 于: 六月 08, 2018, 05:01:49 pm »
本章节主要说明了, 如何计算理论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和最佳性能的,是最常见的方式。但有的时候,这种方式不能用。为何?
[list=1]
  • 很多老板给员工配置的机器,和实际部署的机器的GPU不同。员工的开发的时候无法针对性的对此GPU进行调节优化。此时这种方式无法使用。
  • 在开发的时候,无法知道项目完成后的一定时间内, 例如3年,未来会应用到的新GPU会如何。此时无法直接在未来(现在还没有发布呢)的GPU上跑profiler的。
  • 现在是给广泛性的用户部署,用户可能有现在无法确定下来的配置情况。这种时候,开发时使用的profiler来现场确定的方式是无法生效的。


*

sisiy

  • *****
  • 121
    • 查看个人资料
(无标题)
« 回复 #2 于: 六月 08, 2018, 05:59:55 pm »
此时需要此文中的另外3种方法。这3种方法各有特色。一种是使用excel电子表格。这个表格做的非常不错。如图,这个表格有3个步骤(左侧用户能改动的地方),例如我这里随意输入了:计算能力6.1(Pascal),只使用L2 cache,每个block需要128个线程,每个线程需要48个寄存器,每个block需要4KB的shared memory,然后这个计算器会立刻告诉我(表格的内容会跟随更新)。当前我这种kernel的资源使用,将只能导致63%的occupancy,而且会告诉我(第二行红色)限制因素是因为我的寄存器使用太多了。然后我可以随意输入一点东西。
例如改成这种资源使用。则本表格内容会在下面自动改变:
现在的限制因素变成了我使用了过小的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个的。这种信息就规避了盲目的去优化一些东西。



*

sisiy

  • *****
  • 121
    • 查看个人资料
(无标题)
« 回复 #3 于: 六月 08, 2018, 09:41:00 pm »
因为没有必要。优化了也没有结果。而这些是手册不说的,本表格是直接的一手信息来源(NV),使用该表格需要知道一些信息,例如kernel使用了多少寄存器,这个信息可以在编译的时候通过ptxas的详细输出得到,手册的后面会提到如何得到这个信息。这也是和之前的用profiler获取实际occupancy的不同,该表格只会自动知道设备的信息(例如计算能力6.1的1080),而不能知道你的kernel函数的信息,因为它只是一个普通的表格(好无辜的感觉),但依然已经比你在没有任何卡的情况下,纯靠蒙的方式要好。
[size=10.5pt]本章节还提供了另外一种方式,叫做通过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个只知道结果,不知所以然。而使用这个文件会知道为何(打开看一下就可以了),适合需要进阶的用户,或者手头完全无卡的用户。
[/size]

*

sisiy

  • *****
  • 121
    • 查看个人资料
(无标题)
« 回复 #4 于: 六月 08, 2018, 10:02:39 pm »
关于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目前没有在手册里具体说明此问题。


只有一句话:


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