Qualcomm_Mobile_OpenCL.pdf 翻译-7 内存性能优化

 

内存优化是最重要也是最有效的OpenCL性能优化技术。大量的应用程序是内存限制而不是计算限制。所以，掌握内存优化的方法是OpenCL优化的基础。在这章中，将会回顾OpenCL的内存模型，然后是最优的实践方法。

 

7.1 在Adreno GPU中的OpenCL内存模型

         OpenCL定义了四种内存类型——也就是，global（全局的），local（本地的），constant（常量的），和private（私有的）内存，理解这些内存的不同点是基本要求。图7-1展示了四种内存概念上的设计图。

 

         图7-1 OpenCL概念上的内存结构

 

         OpenCL标准只在概念上定义了这些内存，至于如何实现是由厂商自己定义的。物理上的位置可能与概念上的位置有所不同。比如，private内存对象可能会被放在片外RAM上，离GPU很远。

 

         表7-1列出了在Adreno GPU上4种内存的定义，以及他们的延迟和物理位置。在Adreno GPU上，local和constant内存被放在了片上RAM上，比片外RAM有更短的延迟。

 

         通常地，对经常要访问到的数据，建议使用local和constant内存，以便能更好地利用这种短延迟的特性。更多的细节将会在接下来的章节中说明。

 

         表7-1 Adreno GPU上的OpenCL内存模型

内存类型	定义	相对延迟性	位置
Local	被一个work group内的所有work item共享	中等	片上，SP内部
Constant	work group内部所有的work item的常量数据	当放到片内是，延迟低。否则，延迟比较大	如果能放下，会放在片内，否则放到系统RAM中。
Private	对于work item单独拥有	由编译器决定将它分配到哪块内存	在SP上放在寄存器或者local内存或者系统RAM(由编译器决定）
Global	可以被所有的work group中的所有work  item访问	大	系统RAM

 

7.1.1 local内存

         Adreno GPUs支持快速的片上local内存，不同系列/等级中的local内存大小都有所不同。在使用本地内存之前，比较好的做法是，先使用下面的API函数，查询一下该设备上每一个workgroup可以使用的本地内存有多大。

 

         clGetDeviceInfo(deviceID, CL_DEVICE_LOCAL_MEM_SIZE, .. )

 

下面是使用本地内存的一些方法：

• n  使用本地存储反复使用的数据，或者存储kernel函数中两个阶段之间的临时结果。
  • 　　一种理想的场景是，work items多次获取相同的内容数据时，而且大于2次。

    比如说，考虑这种情况，在某种视频处理中，使用对象匹配方法的滑窗方法。假设每个work item使用带16x16像素的滑窗处理8x8的像素区域，那么相邻的work item之间有大量的数据重叠。在这种情况下，本地内存很适合存储像素，减少了多余的数据获取。

 

•  在不同的work items之间使用barrier进行同步很可能非常耗时
  • 　　如果存在work item之间的数据交换，比如说，work item A向本地内存中写数据，work item B从本地内存中读取数据，因为OpenCL的内存一致性模型并不严格，所以需要一个barrier操作。
  • 　　Barrier 经常会导致同步延迟，从而阻塞ALU，导致更低的ALU的使用效率。
  • 　　在某些情况下，将数据缓冲到本地内存中可能会需要同步，同步产生的延迟将会抵消使用本地内存带来的性能提升。在这种情况下，直接使用全局内存，避免使用barrier可能是更好的选择。
• 使用向量化的装载和存储本地内存
  • 　　可以使用高达128位的向量装载方式，建议装载时32位对齐。
  • 　　在7.2.2节中，将会对向量化的存储和装载做更多细节地讨论。

 

• 让每个work item 参与本地内存的数据装载，而不是用一个work item来完成整个装载任务。
  • 　　避免使用一个work item 为整个work group装载或者存储所有的本地内存。
• 避免调用函数async_work_group_copy。对编译器来说，自动生成装载本地内存的最优的代码是非常困难的，所以开发者通过代码手动装载数据到本地内存是一种更好的方式。

 

7.1.2 constant 内存

         Adreno GPU支持片上的constant 内存。这种类型的内存在4种类型的内存中，有最好的延迟性和更高的性能。constant内存一般用在以下情况中：

• 使用constant修饰符定义的常量或者向量将会被存储在constant RAM中。
• 使用constant修饰符定义的数组会被存储在constant RAM中，如果数组是在程序范围内定义（比如编译器可以决定它的大小）而且在constant RAM上有足够的空间。
• kernel参数中的常量或者向量数据将会被存储在constant RAM中。比如，在下面的例子中，coeffs将会被存储到constant RAM中

         __kernel void myFastKernel(__global float* bar, float8 coeffs)

　　　　{  //coeffs will be mapped to constant RAM  }

 

• 使用__constant修饰的常量和向量变量以及数组，不适合放到constant RAM中，将会被放到系统RAM中。
• 如果要将kernel函数参数中定义的数组装载到constant RAM中，必须提供一个叫做max_constant_size(N)的属性来表示constant array的大小，其中N表示需要的字节数（bytes）。在接下来的例子中，constant RAM中的1024个字节将会被分配给foo：

         __kernel void myFastKernel( __constant float foo* __attribute__( (max_constant_size(1024)))

         {  . . .  }

 

         指定max_constant_size属性是非常重要的。如果没有这个属性，这个数组将会被存储到片外的系统RAM上，因为编译器并不知道数组的大小，从而不能将它放到片上RAM上。

 

注意：这个特性只能支持16位和32位数据的数组，8位数据的数组并不支持。另外，如果在constant memeory中没有足够的位置分配给数组，那么数组将会被存储到片外的系统RAM上。

 

注意：对于动态索引的数组和被不同的work item访问的数组， constant RAM并不是最优的。比如，一个work item获取索引0，接下来的另一个work item获取索引20，这种constant 内存是无效的。在这种情况下，使用image对象是一个更好的选择。

 

 

7.1.3 private 内存

         在OpenCL中，private内存是每一个work item私有的，不能被workgroup中的其他work item访问。物理上，private内存可以存在在片上寄存器或者片外系统RAM上。它的实际存放位置依赖某些因素，下面是一些典型的情况：

• 常量将会被存放在寄存器中，寄存器是最快的内存。
  • 　　如果没有足够的寄存器，私有变量将会被放到系统RAM中。
• 私有的数组将会被存储在：
  • 　　本地内存，但是这不能保证
  • 　　片外系统RAM，如果数组超出了本地内存的容量

         将private内存存储在片外系统RAM上是非常不理想的，因为系统RAM比较慢，而且私有内存的访问方式并不能很好地使用cache，特别是当每个work item的private内存数量很大的时候。下面是一些建议：

• 避免在kernel中定义任何private 数组。如果可能的话，尝试使用vector。
• 用global或者local内存替代private数组，并设计成当多个work item获取数组元素时可以合并获取。
• 使用向量化进行private 内存装载/存储，比如，在每次处理中尝试装载/存储高达128位的数据。

      7.1.4 global 内存

         OpenCL 应用程序能够使用两种类型的全局内存对象，buffer和image，这两种全局内存对象都是使用片外系统内存。buffer对象是一个简单的一维数组，image对象是一种模糊的内存对象，开发者不能假设数据在内存中存储的布局和格式。当一个image对象创建时，软件会将数据安排成GPU能够更有效的访问形式。使用它们的最有效方式是不同的，这个将会在接下来的章节中讨论。

       7.1.4.1 Buffer

         Buffer对象存储一维元素的集合，这些元素可以是数值类型的数据（比如整型，浮点型），向量数据类型或者用户定义的数据结构。一个buffer对象可以使用以下的API函数创建：

         cl_mem clCreateBuffer (cl_context context,

                            cl_mem_flags flags,

                            size_t size,

                            void *host_ptr,

                            cl_int *errcode_ret)

 

         Buffer 对象存储在global内存中，而且在Adreno GPUs中可以通过L2 cache访问。在这个函数中，最重要的参数是cl_mem_flags。OpenCL允许这个函数有很多不同的标志，如果选择和结合这些标志对性能提升是非常重要的。下面是一些建议：

• 一些标志会导致额外的内存拷贝。尝试使用zero-copy 标志，这个标志将会在7.4节中讨论。
• 一些标志是针对台式/分离式的GPU，这些GPUs有自己专属的GPU内存。
• 使用最精确的标志。总的来说，标志越严格，OpenCL的驱动就能找到越是适合的配置，从而提高性能。比如说，它可以找到最适合该内存对象的cache刷新的规则（写过去，写回来等）。7.4.2节中有对cache规则和对性能影响的详细说明。下面是一些例子：
  • 　　如果内存只是被host端读，使用CL_MEM_HOST_READ_ONLY
  • 　　如果内存无法被host端访问，使用CL_MEM_HOST_NO_ACCESS
  • 　　如果内存只能被host端写，使用CL_MEM_HOST_WRITE_ONLY

 

7.1.4.2 Image

         一个Image对象是用来存储一维，二维，或者三维的纹理，帧缓冲，或者一个图像数据，image对象中的数据布局是不透明的。实际上，这个对象中的内容并不需要必须与一个真实的图像数据有关系。任何数据都可以存储成image对象的格式，这样就可以在Adreno上使用硬件的texture引擎和它的L1 cache。

         一个image对象可以使用下面的API创建：

         cl_mem clCreateImage(cl_context context,

                         cl_mem_flags flags,

                         const cl_image_format *image_format,

                         const cl_image_desc *image_desc,

                         void *host_ptr,

                         cl_int *errcode_ret)

 

         注意，image对象的cl_mem_flags与前面讨论的buffer对象中的标志有相同的规则。

         Adreno GPUs支持很多种图像格式和数据类型。从Adreno A3x GPU到Adreno A5x GPU，添加了新的image格式和数据类型。用户可以使用函数clGetSupportedImageFormats去获取支持的图像格式和数据类型的完整列表。

 

         为了充分使用内存带宽，建议使用数据长度是128位的标志对，比如CL_RGBA/CL_FlOAT, CL_RGBA/CL_SIGNED_INT32等。

 

       7.1.4.3 使用image对象使用vs. buffer对象

         相比buffer对象，image对象有以下优点：

• 能使用texture 引擎硬件。
• 使用L1 cache
• 内嵌有图像边界的处理。
• 支持大量的图像格式和数据类型的结合，已经在7.1.4节“Image”列出，同时支持自动的格式转换。

 

         OpenCL支持两种简单的滤波，CLK_FILTER_NEAREST 和 CLK_FILTER_LINEAR。对于CLK_FILTER_LINEAR，结合适合的image类型，能够让GPU使用内嵌的texture 引擎做自动的双线性差值。

         举例来说，假设一个图像的类型是CLK_NORMALIZED_COORDS_TRUE和CL_UNORM_INT16，假设图像数据是2字节的unsigned short类型。read_imagef的函数调用将会做以下的工作：

• 从image对象中读取像素点（这些像素点将被缓冲到L1 cache中）    
• 在硬件上进行临近像素点差值。
• 转换并归一化到[0,1]。

 

         对于双线性或者三线性差值操作来说，这个很方便。

         有时，buffer对象可能是个更好的选择：

• 更灵活的数据获取方式：
  • 　　image对象只能允许按像素大小的边界访问，比如，对于RGBA的128位，32位/通道

• • 　　对于buffer对象，Adreno支持字节寻址访问。比如，在buffer对象中，在没有超过buffer边界的情况下，128位数据可以从任何字节地址装载

•  如果L1是瓶颈
• 比如说，出现很严重的L1 cache垃圾，使得L1 cache访问效率很低。
• 一个buffer对象可以在kernel中读和写。尽管image对象从OpenCL2.0开始，也能够读和写，但是由于同步的要求，它的性能很低。

 

         表7-2 Adreno GPU上的Buffer vs. image

        

特性	Buffer	Image
L2 Cache	Yes	Yes
L1 Cache	No	Yes
支持对象的读和写	Yes	No（在OpenCL1.x） Yes （在OpenCL2.x，有同步的要求）
按字节寻址	Yes	No
带内嵌的硬件插值	No	Yes
带内嵌的边界处理	No	Yes
支持Image格式和采样	No	Yes

 

 

7.1.4.4 同时使用Image和buffer对象

         相比于仅使用texture对象或者buffer对象，一个更好的方式是同时能够充分使用UCHEó SP 和UCHEóTPL1óSP 这种路径。因为TPL1有L1cache，将最常用的但是相对数量少的数据存储在L1上是一个好的方法。

 

7.1.4.5 Global 内存 vs. Local 内存

         local内存的一种使用法是，先将数据装载到本地内存，进行数据同步，保证数据已经可用，然后work group的work item用这些数据进行处理。但是，可能由于以下几个原因，使用全局内存可能比本地内存更好：

•  可能有更好的L2 cache的命中率和更好的性能。
• 代码会比使用本地内存时简单，而且会有更大的work group尺寸。

        

      7.2 优化内存的装载/存储

         在之前的章节中，我们讨论了如何使用不同类型的内存。在这节中，我们将仔细考察，内存的装载/存储对性能影响的一些关键的和普遍的重点。

      7.2.1 合并的内存装载/存储

         合并装载/存储指的是把多个相邻的work item装载/存储的请求合并的能力，这个已经在3.3.1节讨论local 内存访问时提到过。合并访问对于global内存的存储/装载也很重要。

 

         除了装载是两路处理（请求和响应），存储是一路处理这种情况以外，合并存储跟读操作的工作方式类似。因此，合并的装载比存储更严格。

 

         在Adreno GPUs中，从Adreno A4x系列的GPU开始，硬件的合并操作逐渐被使能，如表7-3所示。Private内存不支持合并访问。

        

         表7-3 在Adreno GPUs中支持合并访问的GPUs系列

装载/存储	Adreno A3x	Adreno A4x	Adreno A5x
global内存合并装载	No	No	No
global内存合并存储	No	Yes	No
本地内存合并装载/存储	No	No	Yes

 

      7.2.2 向量化的装载/存储

         向量化的装载/存储指的是一个work item使用向量化的方法同时装载/存储多个数据。这个与合并访问是不同的，合并访问时是多个work item。下面是使用向量化装载/存储的一些关键点：

• 对于每一个work item，建议同时装载一整块数据，比如64bit/128bit，这样能够更好地利用带宽。
  • 　　比如，多个8位的数据可以手动打包成一个元素（比如64位或128位），这样可以使用vloadn装载，然后通过as_typeN函数（比如as_char16）进行解包。
  • 　　可参考9.2.3中向量化操作的列子。

 

• 为了更好的优化SP到L2的带宽性能，装载/存储的内存地址必须是32字节对齐的。
• 有两种方法来向量化装载/存储
  • 　　使用内嵌的函数（vloadn/vstoren），这些函数已经在OpenCL中被很好的定义了。
  • 　　另外，指针的强制类型转换同样也可以用来向量化的装载/存储，如下所示：

                            char *p1; char4 vec;

                            vec = *(char4 *)(p1 + offset);

•  建议在使用向量化的装载/存储指令时，最多使用4个元素（vload4/store4）。因为当向量化装载超过4个元素的数据类型时，将会被拆分成多个装载/存储指令，这些指令的操作数不会超过4个元素

• n  避免在一个work item中装载太多的数据

• 装载太多的数据可能会使用更多的寄存器，从而导致更小的工作组尺寸，以及性能的损失。在最坏的情况下，会引起寄存器溢出，比如，编译器可能需要使用系统RAM来存储变量。

注意：向量化的ALU计算同样也可以提高性能，尽管一般来说，没有向量化的内存/存储对性能提升的多。

 

7.2.3 优化数据类型

         数据类型也很重要，因为它不仅影响内存的繁忙程度，也会影响ALU的操作。下面是数据类型的一些规则：

•         检查程序流中的每一个阶段的数据类型，保证在整个流程中每个阶段的数据类型一致。
•        如果可能，使用更短的数据类型，这样可以减少内存的获取/带宽，而且可以提高可执行的ALU数量。

7.2.4 16位宽的浮点 vs. 32位宽的浮点

         因为Adreno GPUs有专用的硬件来计算half-float数据类型，所以强烈建议使用halt-float来替代float数据类型。half ALU的gflops几乎是full ALUs的两倍。下面是一些规则：

• 16位宽的half的支持精度是有限的。它仅能够精确的表示很小范围的数据。
  • 　　比如，它只能够精确地表示在[0,2048]整数范围内的数据。
• 如果half的数据计算会导致不能接受的精度损失，那需要将half转成float精度。但是在存储时，仍然使用half数据类型。

7.3 原子函数

         在OpenCL中定义了大量的本地的和全局的原子函数，而且Adreno GPU在硬件上就支持。下面是使用原子函数的一些规则：

•         避免一个或者多个工作组频繁去更新一个单独全局原子内存地址，因为原子操作是顺序操作，而且他们的性能比如并行操作。
•       首先应该尝试使用本地原子操作，然后再对全局内存进行一次原子更新。

 

7.4  0拷贝

         Andreno OpenCL提供了一些机制来避免可能在host端发生耗时的内存拷贝。因为与内存对象创建的方式有关，所以有一些不同的方法来避免耗时的拷贝。

7.4.1 使用map替换拷贝

         假设OpenCL应用程序对数据流有完全的控制权，比如目标和源内存对象的创建都是由OpenCL应用程序管理的。这是一种最简单的情况，可以使用以下几个步骤避免内存拷贝：

• 当创建一个buffer/image对象时，使用CL_MEM_ALLOC_HOST_PTR标志，然后执行以下的步骤：
  • 　　首先在clCreateBuffer函数中设置cl_mem_flags：

                  cl_mem Buffer = clCreateBuffer(context,

                         CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,

                         sizeof(cl_ushort) * size,

                         NULL,

                         &status);

 

• • 　　然后使用map函数，给host返回一个指针

                   cl_uchar *hostPtr = (cl_uchar *)clEnqueueMapBuffer(

                                commandQueue,

                                Buffer,

                                CL_TRUE,

                                CL_MAP_WRITE,

                                0,

                                sizeof(cl_uchar) * size,

                                0, NULL, NULL, &status);

 

 

• • 　　host使用指针hostPtr来更新buffer

          比如，host可以读取相机数据或者从硬盘中的数据填充到buffer中。

• • 　　取消对象映射

                                status = clEnqueueUnmapMemObject(

                     commandQueue,

                     Buffer,

                     (void *) hostPtr,

                     0, NULL, NULL);

 

 

• • 　　然后这个对象可以被OpenCL的kernel函数使用

 

         在这种场景下，CL_MEM_ALLOC_HOST_PTR是避免数据拷贝的唯一方法。使用其他的标志，比如CL_MEM_USE_HOST_PTR和CL_MEM_COPY_HOST_PTR，为了GPU能访问这些数据，驱动将必须要进行额外的内存拷贝

      7.4.2 避免不是由OpenCL分配的对象的内存拷贝

       7.4.2.1 ION 内存扩展

         如果内存对象是在OpenCL API范围外被初始化创建的，而且是使用ION/Gralloc分配的，那么可以使用 cl_qcom_ion_host_ptr这个扩展来创建buffer/image对象，这样会将ION内存映射到GPU可访问的内存，而且不会发生额外的拷贝。

 

注意：如果有需要，可以提供一个详细的简单的代码来阐述通过QTI扩展来使用ION内存而避免内存拷贝的方法。

 

       7.4.2.2 QTI ANB(Android native buffer）扩展

         在许多相机和音频处理的使用案例中，ANB（由gralloc分配的）必须共享。因为buffer是基于ION的，所以共享是有可能的。然而，为了使用ION，开发者需要从从buffer中获取内部的句柄，这个需要访问QTI的内部头文件。cl_qcom_android_native_buffer_host_ptr扩展提供了一个更直接的方式与OpenCL共享ANB，并且不需要访问QTI的头文件。这样ISV和其他第三方的开发者能够在ANB上实现0拷贝。

        

注意：如果有需要，可以提供一个例子来阐述 cl_qcom_android_native_buffer_host_ptr扩展的使用。

        

       7.4.3 使用标准的EGL 扩展

         cl_khr_egl_image 扩展能够从EGL图像中创建一个OpenCL image对象。这样做的主要好处是：

• 这是一种标准，使用这种技术书写的代码能够最大程度上在其他支持EGL的GPU上工作。
• 由于EGL/CL 扩展（l_khr_egl_event 和EGL_KHR_cl_event）的设计，所以使用EGL/CL 扩展的程序可能会实现更有效的并行。
•  对于YUV处理，使用EGL_IMG_image_plane_attribs扩展会更简单。

 

7.5 提高cache的利用

         为了有更好的利用cache，必须要遵守下面的一些规则：

 

• 检查cache的垃圾和cache的使用率。Snapdragon Profiler可以提供cache访问的信息，比如装载/存储的字节数，cache命中率/没命中率。
  • 　　如果装载进UCHE的字节数比kernel期望的多很多，那么可能存在cache垃圾。
  • 　　比如L1/L2 命中/没命中率等指标能够提供cache的使用情况。

 

• 通过以下方法避免cache垃圾
  • 　　调整workgroup大小，比如减少workgroup的大小。
  • 　　改变访问模式，比如，改变kernel的维度。
  • 　　如果使用loops时产生cache垃圾，在循环中添加automics或者barrier可能会减少垃圾。

                  

7.6 CPU的cache操作

         对于可以用cahce缓存的内存对象，OpenCL驱动需要在合适的时间更新cache数据或者使cache数据无效。这能够保证当CPU和GPU尝试访问数据时，他们看到的是最新的数据拷贝。比如，当主CPU为了读数据，映射一个kernel的输出buffer时，那必须使CPU cache中的数据无效。OpenCL的驱动程序有非常复杂的CPU cache管理机制，通过对每一块内存对象偏移的可视化跟踪和尽可能推迟额外的操作，尝试使用最少数量的cache操作。比如，在一个kernel启动前，可能会对输入buffer的CPU cache的进行刷新。

         CPU cache操作会有可以测量的损失，可以通过clEnqueueNDRangeKernel中的CL_PROFILING_COMMAND_QUEUED与CL_PROFILING_COMMAND_SUBMIT之间的差值查看，如图4-1显示的那样。在某些情况下，clEnqueueMapBuffer/Image和clEnqueueUnmapBuffer/Image的执行时间可能会增加。一个CPU cache操作的耗时通常会随着内存对象的大小线性增加。

         为了最小化CPU cache操作的耗时，必须对应用程序的流程进行仔细地安排，避免在CPU和GPU之间来回切换处理。而且，应用程序分配内存对象时，在CPU和GPU之间来回被访问的数据和只有一种访问的数据需要放在不同的内存对象里。

        

         内存对象创建时必须使用CPU cache机制，这个机制需要跟他们的用途合适。当为buffer对象或者image对象分配内存时，驱动将会选择CPU cache机制。默认的CPU cache机制是write-back。然而，如果使用了CL_MEM_HOST_WRITE_ONLY或CL_MEM_READ_ONLY标志，驱动将会认为应用程序不准备使用host CPU来读取数据。在这种情况下，CPU的cache机制被设置为write-combine。

 

         对于外部的分配内存对象，比如使用ION和ANB机制，应用程序对CPU cache机制更直接的控制。将这些对象引入到OpenCL时，应用程序必须要正确设置CPU cache机制。

 

7.7 使用SVM

         Adreno A5x GPUs支持粗粒度的SVM，这是OpenCL2.0 完整版简介中一个关键特性。

使用SVM时，host和devie的内存地址是相同。在OpenCL2.0中的SVM特性能够方便地实现host端和device端之间的内存共享，能通过在OpenCL设备上访问host指针。

        

         对于粗粒度的SVM，在同步时（map/unmap）,host或者devices上访问内存将会被限制了.对于需要在host端和device端都进行处理的数据结构指针这类应用程序可以很好的利用这个特性。

 

7.8 减少电源/热量消耗的最好的经验

         对于移动的应用程序，能源和热量是一个主要的考虑因素。高性能的应用程序可能没有最好的电源/热量的性能，而且反之亦然。因此，理解电源/热量和性能的要求是很重要的。下面是一些在OpenCL上减少电源和热量消耗的几个提示：

•  使用所有的方法避免内存拷贝。比如，使用ION 内存来实现0拷贝，而且，在使用函数clCreateBuffer创建buffers使用标志CL_MEM_ALLOC_HOST_PTR。另外，避免使用OpenCL APIs进行数据拷贝。
• 最小化host和device之间的内存传输。可以通过以下方式实现，在constant内存或者local内存存储内存数据，使用更短的数据结构，降低数据精度，剔除private内存的使用等。
• 优化kernel和提高他们的性能。kernel运行的越快，消耗的能量和电量越少。
•  最小化软件的开销。比如，使用事件驱动的流程来减少host和device的通信开销。避免创建太多的对象，避免在kernel执行之间创建或者释放对象。

 

转载于:https://www.cnblogs.com/xiajingwang/p/11120561.html