Chart 7 内存优化

文章目录

前言
7.1 Adreno GPU OpenCL内存
- 7.1.1 内存声明周期
- 7.1.2 Loacl Memory
- 7.1.3 Constant memory(常量内存)
- 7.1.4 Private Memory
- 7.1.5 Global Memory
- - 7.1.5.1 Buffer Object
  - 7.1.5.2 Image Object
  - 7.1.5.3 Image object vs. buffer object
  - 7.1.5.4 Use of both Image and buffer objects
  - 7.1.5.5 Global memory vs. local memory
7.2 优化内存 load/store
- 7.2.1 Coalesced memory load/store
- 7.2.2 Vectorized load/store
- 7.2.3 Optimal data type
- 7.2.4 16-bit vs. 32-bit data type
7.3 Atomic functions in OpenCL 1.x
7.4 Zero copy
- 7.4.1 使用内存映射而不是 COPY
- 7.4.2 避免为非由OpenCL分配的对象进行内存复制
- - 7.4.2.1 ION/dmabuf memory extensions
  - 7.4.2.2 QTI Android native buffer (ANB) extension
  - 7.4.2.3 Android Hardware Buffer (AHB) extension
  - 7.4.2.4 Using standard EGL extensions
7.5 Shared virtual memory (SVM）
7.6 Improve the GPU’s L1/L2 cache usage
7.7 CPU cache operations
7.8 减少能耗
总结

前言

内存优化是最关键且有效的OpenCL性能技术。许多应用程序受限于内存而非计算能力。因此，精通内存优化对于OpenCL优化至关重要。

7.1 Adreno GPU OpenCL内存

OpenCL定义了四种类型的内存（全局、本地、常量和私有），了解它们之间的差异对性能优化至关重要。图7-1说明了这四种内存类型的概念布局。
在这里插入图片描述

OpenCL标准仅在概念上定义了这些内存类型，它们的实现是供应商特定的。物理位置可能与其概念位置不同。例如，私有内存对象可能位于离GPU很远的片外系统内存中。

表7-1列出了Adreno GPU中四种内存类型的定义，以及它们的延迟和物理位置。在Adreno GPU上，本地内存和常量内存都位于芯片上，其延迟比片外系统内存要短得多。

一般来说，内核应该使用本地(Local)和常量内存来存储需要频繁访问以利用低延迟特性的数据。更多详细信息在接下来的章节中会有介绍。
在这里插入图片描述

7.1.1 内存声明周期

一个典型的问题是如何将内存对象的内容从一个内核传递到下一个内核。例如，如何在内核的本地内存中共享内容，以便在后续内核中使用。以下是开发者应该遵循的原则：

本地内存是每个工作组独有的，其内容的生命周期在工作组执行完成后结束。因此，无法在一个工作组的本地内存内容或从一个内核到另一个内核中进行共享。
常量内存内容在工作组中的所有工作项之间是一致的。一旦内核执行完成，内容可能被 GPU 上运行的其他任务（如图形工作负载）覆盖。
单个工作项拥有私有内存，一旦工作项执行完成，就不能共享。
全局内存由主机创建的缓冲区和图像对象支持，可以由主机和 GPU 访问。因此，如果对象没有被释放，它可以通过不同的内核访问。

7.1.2 Loacl Memory

Adreno GPU支持快速的片上 Loacl Memory，但 Loacl Memory 的大小在不同系列/层级 GPU 之间会有所变化。在使用 Loacl Memory 之前，最好使用以下API查询设备每个工作组可用的 Loacl Memory ：

clGetDeviceInfo(deviceID, CL_DEVICE_LOCAL_MEM_SIZE, ... )

以下是使用本地内存的指南： ps: 工作项（work item）是任务的最小执行单元

使用本地内存来存储在内核中的两个阶段（两次操作）之间需要重复访问的数据或中间结果。
- 理想情况是当工作项多次访问相同内容且超过两次时。
  - 例如，考虑使用对象匹配进行视频处理的基于窗口的运动估计。假设每个工作项处理一个16x16像素的搜索窗口内的8x8像素小区域，导致相邻工作项之间存在数据重叠。在这种情况下，使用本地内存可以很好地存储像素，以减少冗余获取。
在工作项之间进行数据同步的屏障可能会很昂贵。
- 如果工作项之间存在数据交换，例如，工作项A将数据写入本地内存，工作项B从中读取，由于OpenCL的松散内存一致性模型，需要进行屏障操作。
- 屏障通常会导致同步延迟，使ALUs停滞，从而降低利用率。
- 在某些情况下，将数据缓存到本地内存会导致同步延迟，抵消了使用本地内存的好处。在这种情况下，直接使用全局内存以避免屏障可能是更好的选择。
使用矢量化的本地内存加载/存储。
- 推荐使用32位对齐的多达128位（例如，vload4_float）的矢量化加载。
- 有关矢量化数据加载/存储的更多详细信息，请参见第7.2.2节。
允许每个工作项参与本地内存数据加载，而不是使用一个工作项执行整个加载。
- 避免仅使用一个工作项来加载/存储整个工作组的本地内存。
避免使用名为async_work_group_copy的函数。对于编译器来说，生成加载本地内存的最佳代码通常很棘手，因此最好由开发人员手动编写代码将数据加载到本地内存中。

7.1.3 Constant memory(常量内存)

Adreno GPUs支持芯片上的常量内存，如果得当使用，可以在四种内存类型中提供卓越的性能。常量内存通常在以下情况下使用：

标量和矢量变量使用 constant 定义。
如果在程序范围内定义了带有 constant 的数组（例如，编译器可以确定其大小），它将适应常量内存。
内核参数是标量或矢量数据类型。例如，以下示例中的coeffs将存储在常量内存中：

__kernel void myFastKernel(__global float* bar, float8 coeffs)
{ //coeffs will be loaded to constant RAM }

标量和矢量变量以及带有 __constant 但不适合常量内存的数组将被分配到系统内存中。

以下是对于常量内存的一个重要建议。如果一个内核具有以下两个特点：

作为内核参数的小数组，例如5x5高斯滤波器的系数。
该数组的元素在子组或工作组内均匀读取。

其性能可以通过使用名为 max_constant_size(N) 的属性将数组加载到常量内存中而显著提高。该属性用于指定为该数组所需的最大字节数。在以下示例中，为变量 foo 在常量内存中分配了1024字节

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

指定 max_constant_size 属性是至关重要的。如果没有这个属性，数组将存储在片外系统内存中，因为编译器不知道缓冲区的大小，无法将其提升到芯片上的常量内存。此功能仅支持16位和32位的数组，即不支持8位数组。此外，如果缓冲区太大而无法适应常量内存，则它将存储在片外系统内存中。

对于动态索引且由工作项发散访问的数组，常量内存可能不是最优选择。例如，如果一个工作项获取索引0，而下一个工作项获取索引20，那么常量内存效率较低。在这种情况下，使用图像对象可能是一个更好的选择。

7.1.4 Private Memory

在OpenCL中，私有内存是每个工作项私有的，其他工作项无法访问。从物理上来说，私有内存可以存在于芯片内寄存器或片外系统内存中。确切的位置取决于多个因素，以下是一些典型的情况：

标量变量存储在寄存器中，这比其他内存更快。
- 如果寄存器不足，私有变量可能存储在系统内存中。
私有数组可能存储在：
- 本地内存中，尽管不能保证。
- 如果超过本地内存容量，可能存储在片外系统内存中。

将私有内存存储到片外系统内存是非常不可取的，原因有两点：

系统内存的延迟远高于寄存器
私有内存访问模式不友好于缓存，尤其是如果每个工作项的私有内存量很大

建议：

避免在内核中定义任何私有数组。尽量使用矢量。
替换私有数组，使用全局或本地数组，并设计其布局，以便在多个相邻的工作项之间合并对数组元素的访问。这样可以改善缓存性能。
使用矢量化的私有内存加载/存储，即尽量每次加载/存储高达128位的数据，使用vload4/vstore4每次加载/存储四个32位元素。

7.1.5 Global Memory

OpenCL支持使用系统RAM的缓冲区（buffer）和图像对象。与缓冲区对象相比，它是在系统RAM中存储的简单一维数据数组，图像对象是一种不透明的内存对象，开发人员无法看到底层数据的存储方式。当创建图像对象时，软件以特定方式安排数据，以便GPU能够高效访问。它们的最佳使用方式是不同的，并在接下来的部分中进行讨论。

7.1.5.1 Buffer Object

缓冲区对象存储一维元素集合：标量数据类型、矢量数据类型或用户定义的结构体。缓冲区对象的内容通过Adreno GPU中的L2缓存由内核加载或写入。可以使用以下API函数创建缓冲区对象：

cl_mem clCreateBuffer(cl_context context
						cl_mem_flags flags,
						size_t size,
						void *host_ptr,
						cl_int *errcode_ret)

在这个函数中，cl_mem_flags 是一个关键的标志，开发人员必须小心使用，因为它可能会显著影响性能。OpenCL允许在这个函数中使用许多不同的标志，对于Adreno GPU，以下是一些关键点：

一些标志可能会导致额外的内存复制。尽量使用第7.4节中描述的零拷贝标志。
一些标志适用于具有专用GPU内存的台式机/独立GPU。

使用最准确的标志：

总体思路是，标志越严格，OpenCL软件越有可能找到对象的最佳配置。
例如，OpenCL软件可以应用最适合内存对象的缓存刷新策略（write-through、write-back等），以在缓存刷新时产生最小的开销。
第7.4.2节详细介绍了缓存策略及其对性能的影响。以下是一些示例：
- 如果内存只能由主机进行读取，则使用 CL_MEM_HOST_READ_ONLY。
- 如果内存对主机没有访问权限，则使用 CL_MEM_HOST_NO_ACCESS。
- 如果内存仅用于主机写入，则使用 CL_MEM_HOST_WRITE_ONLY

7.1.5.2 Image Object

图像对象存储1D、2D或3D纹理、frame buffer 或图像数据，图像对象内部的数据布局是不透明的。在实际应用中，对象中的内容不一定与实际图像数据相关联。任何数据都可以存储为图像对象以利用Adreno中的硬件纹理引擎及其L1缓存。使用以下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)

请注意，图像的 cl_mem_flags 具有与前一节讨论的缓冲区对象类似的经验法则。

Adreno GPU支持许多图像格式和数据类型。随着版本迭代，又增加了更多的图像格式和数据类型。开发人员可以使用函数 clGetSupportedImageFormats 获取设备上可用的完整图像格式/数据类型列表。为了充分利用内存带宽，开发人员应该使用长度为128位的配对，例如 CL_RGBA/CL_FLOAT、CL_RGBA/CL_SIGNED_INT32 等。

Adreno GPU还通过供应商扩展支持OpenCL标准中没有的格式，比如YUV和压缩格式。除了新格式外，许多新函数也对图像对象进行了硬件加速，例如 box filtering、SAD和SSD。更多详情，请参考第9章

7.1.5.3 Image object vs. buffer object

如第6.2节所述，由于具有强大的纹理引擎、专用L1缓存和自动处理越界访问(这有时候并不是一件好事)等多种优势，Adreno GPU在处理图像对象时比缓冲对象表现更好。Adreno GPU支持许多图像格式和数据类型的组合，并且能够进行自动格式转换。

OpenCL支持两种采样器滤波器，即CLK_FILTER_NEAREST和CLK_FILTER_LINEAR。对于CLK_FILTER_LINEAR，适当的图像类型组合允许GPU使用其内置的纹理引擎进行自动双线性/三线性插值。例如，假设一个图像是CLK_NORMALIZED_COORDS_TRUE和CL_UNORM_INT16，即图像数据为2字节无符号短整型。要执行read_imagef，Adreno GPU执行以下操作：

图像对象通过L1缓存读取像素。
执行所有所需像素的插值。
将其转换并标准化到[0, 1]的范围内。

这对于双线性/三线性插值操作非常方便。Adreno GPU还通过供应商扩展支持双三次插值。有关更多详细信息，请参考第9.3.4节。然而，有时缓冲对象可能是更好的选择：

缓冲对象允许更灵活的数据访问：
- 图像对象只允许在像素大小的粒度进行访问，例如，128位的 32-bit/channel 的 RGBA 图像。
- Adreno支持对缓冲对象进行灵活访问，其中指针提供了在数据访问方面极大的灵活性。
L1 缓存成为瓶颈
- 例如，严重的 L1 缓存抖动，使得 L1 缓存访问变得低效
缓冲对象允许在内核中进行读写操作。在内核中支持带有__read_write限定符的读写图像（即图像对象）。由于同步要求，一些旧一代的Adreno GPU上，读写图像的性能可能不如缓冲区好。

7.1.5.4 Use of both Image and buffer objects

充分利用 L2 cache<=>SP 和 L2 cache<=>TPL1<=>SP 两种方法是更好的方法，由于TPL1具有L1缓存，将最常用但相对较小的数据存储在L1缓存中是一个良好的做法。

7.1.5.5 Global memory vs. local memory

一种典型的本地内存使用情况是首先将数据加载到本地内存中，同步以确保数据准备就绪，然后工作组中的工作项可以使用它进行处理。然而，使用全局内存可能比使用本地内存更好，原因如下：

它可能具有更好的L2缓存命中率和更好的性能。
代码比使用本地内存更简单，并且具有更大的工作组大小（local memory 只对应于一个 workgroup）。

7.2 优化内存 load/store

在前面的部分中，我们讨论了如何使用不同类型的内存的一般指导原则。在本节中，我们将回顾一些关于内存加载/存储对性能至关重要的关键要点。

7.2.1 Coalesced memory load/store

合并加载/存储指的是从多个相邻的工作项中合并加载/存储请求的能力，如第3.2.4节中所述，用于本地内存访问。合并访问对于全局内存的加载/存储同样至关重要。

合并存储的工作方式类似于读取，只是加载是一个双向过程（请求和响应），而存储是一个单向过程，通常不会阻塞内核执行。对于大多数用例来说，数据加载远远大于数据存储。因此，合并加载通常比存储更为关键。

Adreno GPU支持对全局内存和本地内存进行合并访问，但不支持对私有内存进行合并访问。

7.2.2 Vectorized load/store

矢量化加载/存储指的是为单个工作项进行的多个数据加载/存储矢量化操作。这与合并访问不同，合并访问是为了各种工作项。以下是使用矢量化加载/存储的一些关键要点：

每个工作项应以多个字节的块加载数据，例如，64/128位。这样可以更好地利用带宽。
- 例如，多个8位数据可以手动打包成一个元素（例如，64位/128位），然后使用vloadn加载，再使用as_typeN函数（例如，as_char16）进行解包。
- 请参见第10.3.3节中的矢量化操作示例。
为了优化 SP 对L2缓存的带宽利用率，加载/存储的内存地址应该是32位对齐的。
有两种方法进行矢量化加载/存储：
- 使用 built-in function (例如vload/vstoren）。
- 或者，可以使用指针转换来进行矢量化加载/存储，如下所示:

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

使用最多四个分量的矢量化加载/存储指令。具有超过四个分量的矢量化数据类型加载将被分成多个加载/存储指令，每个指令最多处理四个分量。
避免一个工作项加载过多的数据
- 加载过多的数据可能导致更高的寄存器占用，从而导致较小的工作组大小，并影响性能。
- 在最坏的情况下，这可能导致寄存器溢出，即编译器必须使用系统RAM来存储变量。

矢量化ALU（算术逻辑单元）计算也可以提高性能，尽管通常不如矢量化内存加载/存储的提升明显。

7.2.3 Optimal data type

数据类型至关重要，它不仅影响内存流量，还影响ALU操作。以下是一些数据类型的规则：

在应用程序流水线的每个阶段检查数据类型，并确保在整个流水线中使用的数据类型是一致的。
如果可能的话，使用较短的数据类型，以减少内存获取（带宽），并增加可用于执行的ALU数量。

7.2.4 16-bit vs. 32-bit data type

在Adreno GPU上强烈推荐使用16位数据类型而不是32位数据类型，原因如下：

16位ALU操作的计算能力（以gflops为单位）是32位操作的两倍，这要归功于Adreno对于16位ALU计算的专用硬件加速逻辑。
与32位数据的加载/存储相比，16位数据的加载/存储可以节省一半的带宽。

特别是对于一些机器学习和图像处理用例，16位浮点数，也称为半浮点（FP16），是非常理想的。请注意，与32位浮点数据（FP32）相比，16位半浮点的数据范围和精度更为受限。例如，它只能准确地表示整数值范围在[0, 2048]内。开发人员必须意识到精度损失的问题。

另一种使用16位的方式是将数据加载/存储为16位，而计算部分可以使用32位，如果精度损失是不可接受的。与使用32位数据相比，这将节省一半的内存流量这是相当不错的提议。

7.3 Atomic functions in OpenCL 1.x

OpenCL 1.x支持本地和全局原子函数，包括atomic_add、atomic_inc、atomic_min、atomic_max等。请注意，此处讨论的原子函数与第7.5节中的共享虚拟内存（SVM）中的原子函数不同。Adreno GPU在硬件上支持所有这些函数。在使用原子函数时，请注意以下一些规则：

避免让多个工作项频繁地对单个全局/本地内存地址执行原子操作。
- 原子操作是串行且不可分割的操作，可能需要在内存地址上进行锁定和解锁。
- 因此，不建议让多个工作项对单个地址进行原子操作。
尽量首先进行归约操作，例如，首先使用本地原子操作，然后以原子方式对全局内存进行单一更新。
- 在Adreno GPU中，每个SP都有自己的本地内存原子引擎。如果使用全局内存原子操作且它们的地址相同，首先执行本地原子操作有助于减少访问冲突。

7.4 Zero copy

Adreno OpenCL提供了一些机制，以避免在主机端可能发生的昂贵内存复制。根据内存对象的创建方式，存在一些选项来防止过多的复制。本节描述了实现零拷贝的一些基本方法，第7.5节介绍了一种更高级的使用共享虚拟内存（SVM）的技术。

7.4.1 使用内存映射而不是 COPY

假设OpenCL应用程序完全控制数据流，即 target 和 source 内存对象的创建都由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 函数返回指向主机的指针：

cl_uchar *hostPtr = (cl_uchar *)clEnqueueMapBuffer( commandQueue,
													Buffer,
													CL_TRUE,
													CL_MAP_WRITE,
													0,
													sizeof(cl_uchar) * size,
													0, NULL, NULL, &status);

主机使用指针 hostPtr 更新缓冲区。
- 例如，主机可以将相机数据填充到缓冲区中，或从磁盘中读取数据到缓冲区

取消映射

status = clEnqueueUnmapMemObject(
								 commandQueue, 
								 Buffer, 
								 (void *) hostPtr,
								 0, NULL, NULL);

OpenCL内核可以使用这个对象。

在这种情况下，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/dmabuf memory extensions

假设一个内存对象最初是在OpenCL API的范围之外创建的，并且是使用 ION/DMA-BUF 进行分配的。在这种情况下，开发人员可以使用 cl_qcom_ion_host_ptr 或 cl_qcom_dmabuf_host_ptr 扩展来创建 buffer / image 对象，这些对象映射到 GPU 可访问的内存，而不需要额外的复制。

ION（Input/Output Memory Management Unit）是Android系统中用于管理内存的一种机制。

DMA-BUF（Direct Memory Access Buffer）是Linux内核中的一种机制，用于在不同设备之间共享内存区域，而无需通过CPU的中介。

7.4.2.2 QTI Android native buffer (ANB) extension

在许多相机和视频处理用例中，由gralloc分配的ANB（Android Native Buffer）必须在多个设备之间共享。由于这些缓冲区基于ION，因此共享是可能的。然而，开发人员需要从这些缓冲区中提取内部句柄以使用ION路径，这需要访问QTI（Qualcomm Technologies, Inc.）的内部头文件。cl_qcom_android_native_buffer_host_ptr 扩展提供了一种更直接的方式，在无需访问 QTI 头文件的情况下与 OpenCL 共享 ANB。这使得独立软件供应商（ISVs）和其他第三方开发人员能够实现对 ANB 进行零拷贝的技术。

7.4.2.3 Android Hardware Buffer (AHB) extension

类似于上面描述的ANB扩展，cl_qcom_android_ahardwarebuffer_host_ptr 扩展提供了一种简单的方法，可以在无需提取内部ION句柄的情况下与OpenCL共享AHB（Android Hardware Buffer），从而实现零拷贝的AHB应用程序。

7.4.2.4 Using standard EGL extensions

cl_khr_egl_image 扩展允许从EGL图像创建OpenCL图像。这带来的主要好处有：

这是一个标准化的方法；使用这种技术编写的代码很可能在支持的其他GPU上也能正常工作。
与此扩展一起使用的EGL/CL扩展（如 cl_khr_egl_event 和 EGL_KHR_cl_event）使更有效的同步变得可能。
使用 EGL_IMG_image_plane_attribs 扩展，对YUV（色度亮度分量）的处理变得更加容易。

7.5 Shared virtual memory (SVM）

作为引入到OpenCL 2.0标准的一个重要和高级功能，SVM（Shared Virtual Memory）允许主机和设备共享和访问相同的内存空间，避免过多的数据复制，例如，现在可以在OpenCL设备上访问主机指针。

SVM有几种类型，GPU可以选择支持。从Adreno A5x GPU开始，支持粗粒度的SVM和更高级的带有原子操作的细粒度缓冲区SVM。

对于粗粒度SVM，内存一致性仅在使用映射/解映射函数（即 clEnqueueSVMMap 和 clEnqueueSVMUnMap）的同步点上得到保证。
- 因此，粗粒度SVM类似于第7.4.1节中描述的零拷贝技术，因为它们都需要映射和解映射操作。
- 尽管如此，粗粒度SVM允许应用程序在主机和设备之间使用和共享基于指针的数据结构。
细粒度缓冲区SVM消除了粗粒度SVM中映射/解映射同步的要求。
- 细粒度缓冲区SVM是一种“无映射”SVM，即主机和设备可以同时修改相同的内存区域。
  - 尽管如此，它仍然需要一定程度的同步。
- 取决于主机和设备之间的数据访问模式，可能需要不同类型的同步。
  - 如果在主机和设备之间对相同数据没有读写依赖关系，例如，主机和设备正在处理SVM内存对象的不同部分，那么就不需要原子操作/栅栏。
    - 在这种情况下，内存一致性在OpenCL同步点得到保证，例如，在调用 clFinish 后，所有数据将是最新的。
  - 如果存在对内存访问顺序的依赖或要求，例如主机修改了某个数据，设备需要使用新数据，则需要使用原子操作或栅栏。
    - 在创建时，SVM缓冲区必须具有标志 CL_MEM_SVM_ATOMICS。
    - 在内核内部，必须使用 memory_scope_all_svm_devices。
    - 必须使用一组类似于C11的原子函数，并使用适当的内存作用域、顺序和原子标志。

开发人员需要仔细权衡SVM的利与弊。作为一项高级功能，为GPU实现SVM通常需要精密的硬件设计。实施所有这些高级数据共享和同步可能存在潜在的成本，这些成本开发人员可能未察觉到。在复杂的实际用例中，使用 SVM 的门槛相对较高。开发人员在使用SVM时应谨慎，尤其是在主机和设备之间存在大量数据依赖关系的情况下。在这种类型的用例中，同步成本可能会削弱共享虚拟内存空间的优势。

7.6 Improve the GPU’s L1/L2 cache usage

为了实现良好的缓存利用，开发人员应该遵循以下规则：

了解数据加载/存储的影响：
- 许多内核从全局内存加载的数据比要存储的数据多得多。因此，通过执行合并加载、矢量化加载、使用图像等方式，提高数据局部性并减少对缓存行的需求是至关重要的。
- 然而，数据存储也可能对性能产生重大影响。
  - 对于数据存储，必须首先从系统内存加载缓存行，进行修改，然后写回。
  - 如果数据存储的局部性较差，例如，数据写入了太多的缓存行，内存系统必须加载多个缓存行进行更新。
  - 合并写对性能至关重要，因为它可以提高局部性并减少内存系统对缓存行的需求。
检查并避免缓存抖动，以提高缓存使用效率。
- 缓存抖动指的是在缓存行完全被使用之前被驱逐，然后必须重新获取。这可能导致严重的性能惩罚。
- Snapdragon Profiler可以提供有关缓存访问的信息，例如加载/存储的字节数和缓存命中/失效比率。
- 如果加载到L2缓存的字节数远高于内核的预期，可能存在缓存抖动。
- L1/L2命中/失效比率等指标可以告诉缓存的使用情况有多好。
- 避免抖动的方法包括：
  - 调整工作组的大小和形状。
  - 更改访问模式，例如，更改内核的维度。
  - 如果在使用循环时存在缓存抖动，可以通过在循环中添加原子操作或屏障来减少缓存抖动的机会。
```
// 在循环中使用原子操作或栅栏同步对共享内存的读写，防止缓存行被切出
for (int i = 0; i < size; ++i) {
    barrier(CLK_GLOBAL_MEM_FENCE);
    atomic_operation(&shared_memory[i]);
    barrier(CLK_GLOBAL_MEM_FENCE);
}
```

性能分析工具依赖于硬件性能计数器来生成有关缓存使用情况的指标。由于性能计数器旨在传达有关硬件的信息，因此派生的指标，如L1/L2缓存命中率，可能会产生非直观的结果。例如，可能会看到% L2命中率是大幅的负值，表明加载到缓存中的数据量超过了请求的量。在这类情况下，程序员应该关注性能指标值在优化之间的相对变化，而不是指标的绝对值。

7.7 CPU cache operations

现代SOC（系统芯片）具有多级缓存，骁龙SOC也不例外。对于开发人员来说，了解SOC中GPU/CPU缓存操作的基础知识是有帮助的。

OpenCL驱动程序必须在适当的时候刷新或使CPU缓存无效，以确保对于可缓存的内存对象，当CPU和GPU尝试访问数据时，它们都看到最新的数据副本。例如，当将内核的输出缓冲区映射到由主机CPU进行读取时，必须使CPU缓存无效。

OpenCL软件具有复杂的CPU缓存管理策略，该策略试图通过在每个内存对象的基础上跟踪数据可见性并尽可能推迟操作来最小化缓存操作的数量。例如，在启动内核之前，可能会对输入缓冲区进行CPU缓存刷新。

CPU缓存操作的成本是可以明确测量的，通常可通过观察clEnqueueNDRangeKernel的CL_PROFILING_COMMAND_QUEUED和CL_PROFILING_COMMAND_SUBMIT之间的时间差来体现，如图4-1所示。在某些情况下，clEnqueueMapBuffer/Image和clEnqueueUnmapBuffer/Image的执行时间可能会增加。总的来说，CPU缓存操作的成本通常随着内存对象的大小呈线性增长。

以下是减小CPU缓存操作成本的一些建议：

应该设计应用程序的结构，以便不频繁地在CPU和GPU之间移动处理。
此外，应用程序应该分配内存对象，以便需要在CPU和GPU之间进行交替访问的数据与仅有一个访问转换的数据位于不同的内存对象中。
内存对象应该使用适用于其预期使用方式的适当CPU缓存策略创建：
- 在为缓冲区或图像对象分配内存时，驱动程序将选择CPU缓存策略。默认的CPU缓存策略是写回(write-back)。
- 但是，如果在标志中指定了CL_MEM_HOST_WRITE_ONLY或CL_MEM_READ_ONLY中的任何一个，驱动程序将假定应用程序不打算使用主机CPU读取数据。在这种情况下，CPU缓存策略被设置为写合并(write-combine)。
对于外部分配的内存对象，如ION和ANB机制（参见第7.4.2节），应用程序对CPU缓存策略具有更多、更直接的控制。
- 当将这些对象导入到OpenCL时，应用程序必须正确设置CPU缓存策略标志。

7.8 减少能耗

电源和能耗是移动应用的重要因素。在性能最佳的应用可能不具备最佳的功耗/能效性能，反之亦然。因此，了解功耗/能耗和性能需求至关重要。以下是几个减少OpenCL功耗和能耗的建议：

尽量避免内存复制，例如，使用ION内存实现零复制，并在使用clCreateBuffer创建缓冲区时使用CL_MEM_ALLOC_HOST_PTR。此外，避免使用进行数据复制的OpenCL API。
最小化主机和设备之间的内存交互，例如，在常量或本地内存中存储数据，使用较短的数据类型，降低数据精度，消除私有内存使用等。
优化内核并提高其性能。通常，内核运行得越快，消耗的能量或功率就越少。
减少软件开销。例如，事件驱动的流水线降低了主机和设备通信的开销。避免创建过多的OpenCL对象，并避免在内核执行之间创建或释放OpenCL对象。

clEnqueueNDRangeKernel
参数解析