Qualcomm_Mobile_OpenCL.pdf 翻译-6-工做组尺寸的性能优化

对于许多kernels来讲，工做组大小的调整会是一种简单有效的方法。这章将会介绍基于工做组大小的基础知识，好比如何获取工做组大小，为何工做组大小很是重要，同时也会讨论关于最优工做组大小的选择和调整的通常方法。

 

6.1 获取最大的工做组尺寸

         在运行完clBuildProgram后，使用下面的API函数能够查询设备的最大工做组尺寸。

       size_t maxWorkGroupSize;

       clGetKernelWorkGroupInfo(myKernel,

                                myDevice,

                                CL_KERNEL_WORK_GROUP_SIZE,

                                sizeof(size_t),

                                &maxWorkGroupSize,

                                NULL );

 

         在clEnqueueNDRangeKernel中使用的实际工做组尺寸不能超过maxWorkGroupSize。若是应用程序没有指定工做组大小，Adreno OpenCL软件可能会选择最大的工做组尺寸。

        

6.2 须要的和优先的工做组尺寸

         每个kernel函数都有他须要或者优先的工做组大小。对于须要的工做组大小，OpenCL经过下面方法提供给编译器。

• 使用reqd_work_group_size属性。

        

         做为需求，reqd_work_group_size(X, Y, Z) 属性会传入一个指定的工做组尺寸。若是指定的工做组大小不能知足将会返回一个错误。

         好比，若是要求16x16的工做组尺寸：

         __kernel __attribute__(( reqd_work_group_size(16, 16, 1) ))

       void myKernel( __global float4 *in, __global float4 *out)

       {  . . .  }

 

• 使用work_group_size_hint 属性

 

         OpenCL会尝试使用这个指定的尺寸，可是不保证真实的大小与指定的一致。好比，提示使用64x64工做组尺寸：

         __kernel __attribute__(( work_group_size_hint (64, 4, 1) ))

       void myKernel( __global float4 *in, __global float4 *out)

       {  . . .  }

 

         在许多状况下，当工做组尺寸严格指定时，编译器不能保证能编译出最优的机器代码。并且，若是片上寄存器不能知足要求的工做组尺寸时，编译器可能会须要将寄存器溢出到系统的RAM内存上。所以，这两种属性并不建议使用，除非必须指定工做组尺寸，kernel才能运行。

         注意：为了交叉编译的兼容性，将kernel写成依赖固定工做组的尺寸或者布局，并非一个好的作好。

6.3 影响工做组最大尺寸的因素

         若是没有指定工做组尺寸的属性，一个kernel的最大工做组尺寸依赖如下的几个因素：

• kernel的寄存器使用。通常来讲，kernel越复杂，寄存器使用越多，支持的最大的工做组尺寸越小。过多地使用寄存器的缘由以下：
  • 　　一个工做项中有过多的工做任务。
  • 　　有控制流
  • 　　高精度的数学函数（好比，没有使用内部函数或者快速数学运算的编译选项-fastmath）
  • 　　本地内存，若是须要分配额外的寄存器暂时存储装载和存储指令中源和目的地址。
  • 　　私有内存，好比为每个工做组定义了一个数组
  • 　　循环展开
  • 　　内联函数

        

• 通用寄存器的大小
  • 　　Adreno低级系列的GPU可能有更少的寄存器数量。

 

• n  kernel中的栅栏（Barrier）
  •         若是一个kernel没有使用栅栏（barrier），在Adreno A4x and A5x系列中，在不用考虑寄存器使用的状况下，工做组最大能够设置为DEVICEMAXIUMUM。

        

      6.4 没有barrier的kernels

         之前地，一个workgroup中全部的work item要求在同一时间同时驻留在GPU上。对于大量消耗寄存器的kernel，这将会限制他们的最大工做组尺寸，并将会远远小于设备支持的最大工做组尺寸。

         从Adreno A4系列起，不须要考虑寄存器的使用状况，没有barrier的kernel就能够有Adreno支持的最大工做组尺寸，通常是1024。对于这种类型的kernel（没有barrier）来讲，由于不须要wave之间进行同步，因此当一个旧的wave执行完毕，新的wave就能够开始执行了。

 

         在某些状况下，拥有最大的workgroup尺寸并不意味着他们有最好的并行性。一个没有barriers的kernel可能会由于太复杂致使只有不多的wave在SP上并行执行，这将会致使性能下降。开发者须要继续优化和减小寄存器使用，不考虑从clGetKernelWorkGroupInfo函数中获取到的最大的workgroup尺寸。

 

      6.5 工做组尺寸的调整

         这个部分将会介绍一些在选择最优的工做组尺寸和形状时通用的指导准则。

        

      6.5.1 避免使用默认的工做组大小

         若是一个kernel调用没有指定workgroup的尺寸，那么OpenCL会用简单的方法找一个能用的工做组尺寸。开发者必需要要意识到，这种默认的工做组尺寸一般不是最优的。有效的作法是，手动尝试使用不一样的工做组大小和维度（2D/3D），而后找出最优的一个。

 

      6.5.2 越大的工做组尺寸，越好的性能？

         对于许多kernel来讲这是正确的，由于增长工做组尺寸可以容许更多的wave运行在SP上，这样可以更好地隐藏延迟和提高SP的使用。

         然而，对于某些kernel来讲，增长工做组尺寸可能会致使性能退化。一种状况是，因为不良的数据局部性和访问模式，越大的工做组尺寸将致使越多的cache垃圾。这个数据局部性的问题在使用texture获取时更加严重，由于texture cache比统一的L2 cache要小。最终，决定最优的工做大小和维度的本质是kernel的数据获取。

 

6.5.3 固定的 vs. 动态的工做组尺寸

         为了避免同设备之间的性能兼容性，避免假设一个工做组尺寸可以适合全部的设备，避免对workgroup尺寸固定编码。一个指定的工做组大小和维度在一个设备上是最优的，在另外一个设备上多是次优的。所以，给定一个kernel，建议针对kernel可以执行的全部设备统计出不一样的workgroup尺寸，而后在运行时对每一个设备选出一个最优的。

 

6.5.1 一维 vs二维 vs三维 workgroup（1D/2D/3D）

         kernel的维度能够会影响性能。取决于work item的数据获取方式，在某些状况下，一个2D的kernel可能会在cache上有更好的数据本地性（数据在cache上），致使更好的内存获取和更好的性能。然而在其余状况下，一个2Dkernel比1D会产生更多的cache 垃圾。建议尝试使用不一样的维度，从而获取最优的性能。

 

6.6 关于workgroup的其余话题

      6.6.1 全局的work size和填充

         OpenCL 1.x 要求一个kernel的全局worksize 必须是workgroup尺寸的倍数。若是应用程序指定的workgroup 尺寸不知足这个条件，那么clEnqueueNDRangeKernel的函数调用将会返回一个错误。在这种状况下，应用程序能够填充全局worksize，保证它是用户指定的workgroup尺寸的倍数。

 

         注意：OpenCL 2.0 取消的这个限制，并且global worksize 并不须要必须是workgroup size的倍数，这种被叫作非归一化的workgroup。

 

         理想状况是，workgroup 尺寸的第一个维度是wave尺寸的倍数（好比说32），这样能充分利用wave的资源。若是不是这种状况，能够考虑填充workgroup的大小来知足这个条件，须要记住，在OpenCL 1.x中，全局的worksize必须填充（保证是workgroup的倍数）。

      6.6.2 残酷地寻找

         由于workgroup 尺寸选择的复杂性，经验经常是发现最优大小和维度的最好方法。

        

         一种选择是，在程序开始时，使用一个与实际的工做任务相同复杂度（可是通常使用比较简单的任务）的唤醒功能的kernel 去动态的寻找最优的workgroup 尺寸。而后将这个选出来的workgroup尺寸用在实际的kernel中。不少商业的标准检查程序就是使用的这种方法。

      6.6.3 在workgroup中避免不均匀的工做负载

         一些应用程序可能被写成，在不一样工做组中出现不均衡的负载。好比说，基于区域的图像处理的用例中会出现一些区域须要比其余区域多不少处理的状况。这种状况须要避免，由于这会致使性能的不可预测性。另外的，若是单个workgroup任务须要太长时间运行的话，会致使上下文切换变的复杂。

         解决这个问题的方法是，使用两个阶段处理策略。第一个阶段可能会收集感兴趣的点和为第二阶段准备数据。工做负载越具备肯定性，在不一样workgroups中进行均等的分配将会更简单。

      6.6.4 工做组的同步

         OpenCL 并不能保证workgroup的执行顺序，并且也没有定义一个工做组同步的机制。不建议有须要依赖工做组顺序的程序。

 

         在实际状况下，可使用atomic 函数或者其余方法，在workgroup之间能够进行有限的同步。好比说，一个应用程序可能分配了一个全局内存对象，这个对象须要被不一样的工做组中的workitem更新。一个workgroup能够管理一个由其余workgroup更新的内存对象。经过这种方式，可能会实现有限的工做组同步。