Chart 4 Adreno OpenCL 应用程序开发

半块方便面

已于 2023-12-06 14:18:01 修改

阅读量149

点赞数

分类专栏： OpenCL 文章标签： c++ gpu算力

于 2023-12-05 14:12:30 首次发布

本文链接：https://blog.csdn.net/qq_35227495/article/details/134805367

版权

OpenCL 专栏收录该内容

11 篇文章 2 订阅

订阅专栏

文章目录

前言
4.1 Android上进行OpenCL应用程序开发
4.2 Adreno OpenCL SDK 和 Adreno OpenCL 机器学习 SDK
4.3 调试工具和技巧
4.4 Snapdragon profiler
4.5 性能分析

前言

本章主要介绍如何 debug Adreno OpenCL应用程序

4.1 Android上进行OpenCL应用程序开发

Adreno GPU 主要在 Android 操作系统和部分 Linux 系统上支持 OpenCL。要开发在 Android 上运行的支持 OpenCL 的应用程序，开发者需要使用 Android 软件开发工具包（SDK）和 Android 原生开发工具包（NDK）。有关 Android SDK 和 NDK，请参阅分别位于 SDK 和 NDK。

在本章和后续章节中，假设开发是在 Android 平台上进行的，并且开发者具有 Android SDK 和 NDK 的经验。在 Linux 上进行应用程序开发应该类似。

在 Snapdragon 平台上进行 OpenCL 开发有一些先决条件：

支持 OpenCL 的 Snapdragon 设备。并非所有 Snapdragon 设备都支持 OpenCL，请参阅表 3-1 获取更多详细信息。
OpenCL 软件。Adreno GPU 上的 OpenCL 依赖于 QTI 专有库。
- 检查设备是否安装了 OpenCL 库。
  - 核心库是 libOpenCL.so，通常位于设备上的 /vendor/lib 目录中。
- 一些供应商可能选择不包含 OpenCL 软件（例如，Google 的 Nexus 和 Pixel 设备）。
OpenCL 必须在 NDK 层运行。
对于开发和测试，不需要 root 访问权限，但在运行 SoCs 时可能需要进入性能模式。

4.2 Adreno OpenCL SDK 和 Adreno OpenCL 机器学习 SDK

开发者可以在 https://developer.qualcomm.com/software/adreno-gpu-sdk/tools 找到最新的 Adreno OpenCL SDK 和 Adreno OpenCL 机器学习 SDK。OpenCL SDK 提供了代码示例和文档，帮助开发者理解并有效地使用最新的 Adreno OpenCL 功能。通过一组专有的 API 函数和手动优化的内核，机器学习 SDK 帮助开发者在 Adreno GPU 上进行机器学习推理和训练应用程序的开发。

4.3 调试工具和技巧

由于 GPU 执行的并行性质，调试 OpenCL 应用程序通常是具有挑战性的。对于内核调试，OpenCL 支持 printf 函数，它类似于 c99 中的标准 printf，但有一些细微的差异。建议通过仅打印必要的变量（使用条件来限制输出）来减轻工作负担，因为printf通常会减慢代码执行速度。例如，可以仅启用有问题的工作组，甚至是单个有问题的工作项（通过在函数 CLEnqueueNDRangeKernel 中设置适当的偏移量）。

了解设备的软件版本是很重要的，因为在较新的版本中可能已经修复了一些错误或问题。要查询软件（驱动程序）和编译器版本，开发者可以使用 API 函数，如 clGetDeviceInfo 或 clGetPlatformInfo。

Adreno GPUs调试技巧

在内核中使用 barrier 或 fence 来防止编译器在其前后重新排序代码。
如果返回错误代码，请查阅 OpenCL 规范获取更多信息。
通过调试单个工作项/像素/工作组/内核来隔离问题。
- 例如，将 global_work_size 设置为 [1]，并将 global_offset 设置为像素坐标 [x]。
如果在 API 函数/内核中观察到崩溃，以下是一些要检查的事项：
- 无效的内存地址。
- 其他 API 函数存在问题。
- 内存没有按预期更新。
- 存在溢出或内存缓冲区大小不正确。
  - 尝试单一像素
- 内核执行尚未完成。
  - 使用 clFinish / clWaitforEvent 确保内核执行完成。
- 结果是否不正确且不稳定。
- 不同的工作项是否写入相同的内存地址，是否存在缺失的同步或屏障。

4.4 Snapdragon profiler

Snapdragon Profiler 是一款免费的性能分析工具，可在 Windows、Mac 和 Linux 平台上运行，允许开发人员分析 Snapdragon 处理器的 CPU、GPU、DSP、内存、功耗、热度和网络数据。它支持 OpenCL 以及许多图形 API，如 OpenGL ES 和 Vulkan。有关详细信息，请参阅snapdragon profiler，您可以在该链接下载适用于 Windows、MacOS 和 Linux 的可执行文件以及用户指南。此外，还可以在 Capturing OpenCL applications in Snapdragon profiler中找到简短的 YouTube 视频介绍。Snapdragon Profiler 提供以下一些针对 OpenCL 性能分析的关键功能。

内核分析器（Kernel Analyzer）：该分析器允许开发人员对给定的内核进行静态分析。它提供了诸如寄存器占用和指令数量等信息，帮助开发人员优化内核。
OpenCL API 跟踪和日志：该分析工具提供给定 OpenCL 应用的 API 跟踪和日志。它使开发人员能够从 API 级别识别和解决瓶颈，并进行应用程序调试。
性能信息：该分析器提供 GPU 忙碌比例、ALU 利用率比例、L1/L2 缓存命中比率等信息，有助于开发人员识别内核中的性能问题。

支持命令行和 Android GUI 应用：该分析工具支持基于命令行的应用程序，同时也支持 Android 图形用户界面应用。

4.4.1 Steps to use SDP

连接到设备
首先通过adb连接到手机，然后点击start session会跳出一个窗口，等待一个配置文件安装完毕即可点击connect连接，通过stf连接的adb会出现一些问题，建议使用数据线连接
设置配置信息
a. 选择layout
b. 确保 Blocking 为true
c. 选择 System Trace Analysis
在命令窗口中启动OpenCL应用程序。
a. 类似于“Pending qxprofiler connection and capture signal…” 的消息应该在启动应用程序的命令窗口中弹出，等待下一步操作。

b. 理想情况下，OpenCL应用程序应该被检测并显示在性能分析器GUI的左侧面板上。
选择要进行性能分析的指标。

a. 点击应用程序名称，然后在左下面板中会出现OpenCL Trace和OpenCL Metrics。
b. 选择用于性能分析的追踪和指标。
收集结果。
a. 点击开始进行捕获，应用程序应该恢复执行直到完成。
b. API函数的历史记录可在主窗口中查看，并可进行缩放。
c. 性能分析指标显示在底部窗口中。

d. 点击“Capture->New trace”可以启动新的追踪会话。
e. 可以将结果导出为CSV文件以进行离线分析。

4.4.2 如何解释SDP（Software Development Platform）中的指标

以下是OpenCL应用程序性能分析的一些关键指标

ALU 利用率 %。
- 低值可能表明内核可能受制于内存。
- 提高数据加载/存储的效率。
L2 全局缓冲读取 %。
- 低值可能表示L2缓存未充分利用（可能发生缓存抖动）。
- 平衡工作组的工作负载并调整工作组大小。
L1 命中率 %。
- 如果内核不使用图像对象，则为0。
GPU 忙碌 % 和空闲 %。
- GPU在执行内核时应该是完全忙碌的。
- 理想情况下，忙碌 % 应该接近100%。
- 低于90%的任何数值都应该是一个警告信号，表示主机中有一些问题导致GPU处于空闲状态，而CPU却繁忙。
  - 使用基于事件的流水线，并减少设备与主机之间的同步。

4.4.3 如何有效使用性能分析工具

识别瓶颈。
- 如果某个指标和性能没有改善，而其他指标有所改善，那么未改变的指标可能是瓶颈。
ALU 受限 vs 内存受限 vs 延迟受限：
- 大多数实际情况是内存受限的。
- 有些可能是 ALU 受限，例如在某些卷积或矩阵乘法情况下。
- 延迟受限意味着可能没有足够的波来隐藏延迟。
  - 有关如何调整工作组大小的更多信息，请参见第6.1节。
将计数器的数据与理论数值进行比较。
- 识别缓存抖动：
  - 开发人员可以检查加载到GPU的数据量。
  - 如果字节数超过理论上所需的数据，很可能会发生缓存抖动。
- 开发人员可以设计一些简单的微基准测试来帮助理解GPU的行为。

4.4.4 SDP: static code analysis

在这里插入图片描述
性能分析工具具有静态代码分析器工具，开发人员可以使用它来获取有关内核的基本信息。以下是其中一些最重要的信息。

所有指令。
ALU（算术逻辑单元）/half ALUs（半个ALU）/EFUs（执行单元）。
缓冲加载/存储。
- LDG/STG：全局缓冲加载/存储。
- LDL/STL：本地内存加载/存储。
- LDP/STP：私有内存加载/存储（这通常表示寄存器溢出）。
完整寄存器和半寄存器。
- 确定波的数量。
- 寄存器越多，活跃波越少。
波的最大数量/最大波。
- 如果最大波 < 4，则优化内核/降低复杂性。
  
  以下是开发人员在所有数据中应该处理的最关键问题。
尽一切可能去除 STP/LDP。
- 有关更多详细信息，请参见第7.1.4节。
优化代码以减少 “total footprint” 并增加 “maximum number of waves.”。
- 这两者高度相关。
- “total footprint” 占用越小，GPU可以并行执行的波就越多，性能就越好。

4.5 性能分析

鉴于一个应用程序，准确地对其性能进行分析是至关重要的。以下各节将讨论两种常用的方法，即CPU计时器和GPU计时器，以及它们的主要区别。

4.5.1 CPU Timer

开发人员应该使用C/C++编程语言标准库中的日期和时间函数来测量从主机端调用OpenCL的完整执行时间。一个示例是使用gettimeofday，如下所示：

#include <time.h>
#include <sys/time.h>

void main () {
	struct timeval start, end;
	gettimeofday(&start, NULL); /*get the start time*/
	/*Execute function of interest*/ { . . .
	clFinish(commandQ);
}
gettimeofday(&end, NULL); /*get the end time*/

/*Print the total execution time*/
printf("%ld\n", ((end.tv_sec * 1000000 + end.tv_usec) - (start.tv_sec * 1000000 + start.tv_usec))); 
}

一些OpenCL运行时API函数在其名称中包含“enqueue”，接受一个标志参数来指示是阻塞还是非阻塞调用。对于非阻塞调用，必须谨慎使用CPU计时器。

非阻塞调用意味着在其提交后（通常在另一个CPU线程中排队等待执行），主机继续执行下一条指令，而不是等待函数调用完成。
- 内核执行API函数 clEnqueueNDRangeKernel 是一个非阻塞函数。
对于非阻塞调用，GPU的执行时间不是函数调用之间的时间差。

当使用CPU计时器从主机端测量内核执行时间时，开发人员必须确保使用 clWaitforEvent 调用（如果对于非阻塞调用存在事件ID）或 clFinish 来确认函数的完成。相同的规则适用于内存复制的API函数。

4.5.2 GPU timer

OpenCL的入队函数调用可以选择向主机返回一个事件对象，OpenCL性能分析API可以使用该对象来查询执行时间。Adreno GPU具有用于测量函数执行流的时钟和计时器。 GPU执行时间由GPU硬件计数器提供，独立于操作系统。

为了启用GPU计时器功能，开发人员必须在当前命令队列的clCreateCommandQueue或clSetCommandQueueProperty的属性参数中设置CL_QUEUE_PROFILING_ENABLE标志。此外，必须向入队函数提供一个事件对象。一旦函数执行完成，开发人员可以使用API函数clGetEventProfilingInfo来获取命令执行的性能分析信息。

对于clEnqueueNDRangeKernel调用，使用clGetEventProfilingInfo函数与四个性能分析参数一起，包括CL_PROFILING_COMMAND_(QUEUED, SUBMIT, START和END)，可以提供Adreno GPU中内核启动延迟和内核执行时间的准确图像，如下图所示。

前两个参数的差异，即CL_PROFILING_COMMAND_(QUEUED和SUBMIT)，提供了有关软件开销和CPU缓存操作开销的概念。OpenCL软件可能首先将内核排队，然后与队列中的其他内核一起提交，例如，当队列中的内核数量足够大时。开发人员可以使用clFlush函数加速提交过程。
CL_PROFILING_COMMAND_(SUBMIT和START)之间的差异可以提供GPU正在处理的其他任务的概念。
在GPU上的实际内核执行时间在CL_PROFILING_COMMAND_(START和END)之间。

开发人员应专注于最小化实际内核执行时间，这相对而言比另外两个计时器更为直接，而这两者通常难以控制。
在这里插入图片描述

OpenCL 2.0中引入的 kernel-enqueue-kernel（参见第3.5.2节）功能引入了一个新的性能分析标志，称为CL_PROFILING_COMMAND_COMPLETE。该标志在由事件标识的命令及该命令在设备上入队的任何子命令执行完成时，返回当前设备时间计数器的纳秒值。

4.5.3 GPU timer vs. CPU timer

开发人员可以使用GPU和CPU计时器进行性能分析。尽管GPU计时器可以准确测量GPU执行时间，但一些硬件操作（例如，缓存刷新）和软件操作（例如，CPU主机与GPU之间的同步）超出了GPU时钟系统的范围。因此，对于内核执行，GPU计时器可能会报告比CPU计时器更好的性能指标。以下是两种推荐的实践方法。

使用GPU计时器测量内核优化。从GPU执行的角度，GPU计时器可以精确地告诉每个优化步骤实现的改进量。
如果OpenCL程序仅是整个应用程序流程的一部分，请使用CPU计时器测量应用程序的端到端性能。

4.5.4 Performance mode

骁龙SoCs（系统芯片）具有先进的动态时钟和电压控制机制，可以在特定场景下自动控制系统以运行在省电模式，以节省电池电量。通常情况下，假设存在高强度的工作负载，系统可能会自动提高时钟频率和电压，将设备推入所谓的性能模式，以提升性能并满足工作负载需求。

在给定的OpenCL应用程序中，如果系统动态更改时钟频率，将很难理解和分析其性能。因此，为了实现分析的一致性和准确性，开发人员应启用性能模式。如果没有启用性能模式设置，序列中的第一个OpenCL内核通常会显示更大的启动延迟和较慢的执行时间。开发人员可以在启动实际GPU工作负载之前使用简单的内核来热身GPU。

OpenCL内核的性能不仅仅取决于GPU，CPU主机上运行的API函数与GPU设备上的内核执行同样重要。为了实现最佳性能，CPU和GPU都应启用性能模式。此外，为了减少来自UI渲染的干扰，建议：

确保应用程序以全屏方式呈现，以防其他活动更新屏幕。
如果是原生应用程序，请确保在Android上未运行SurfaceFlinger。这确保应用程序仅由CPU和GPU进行性能分析。
启用性能模式所需的命令序列对于Adreno GPU的不同系列略有不同。更多详细信息请参考附录 A。

4.5.5 GPU frequency controls

应用程序可以利用cl_qcom_perf_hint扩展来控制GPU频率。该扩展允许应用程序在创建OpenCL上下文时设置性能提示属性。性能级别可以是HIGH（高）、NORMAL（正常）和LOW（低）。NORMAL性能级别保留了动态时钟和电压控制的启用。HIGH和LOW性能级别分别禁用了动态电源和时钟控制，并强制GPU以其最大和最小频率运行。

注意：性能级别仅仅是一个提示。驱动程序会尽力尊重这些提示，但诸如热控制、外部应用程序或服务等因素可能会覆盖这些提示。性能提示扩展为应用程序在功耗/性能权衡方面提供了一些灵活性。然而，开发人员应谨慎使用，因为它对SOC级别的功耗有重要影响。

完整的学习文档，在文档读完后上传