Snapdragon上的OpenCL介绍 (3)
4 Adreno OpenCL应用开发
本章简要讨论了Adreno OpenCL应用程序开发的一些基本要求以及如何调试和配置应用程序。
4.1 Android上的OpenCL应用程序开发
目前,Adreno gpu主要支持OpenCL在Android操作系统(OS)和选择Linux系统。要开发一个使用OpenCL运行的Android应用程序,开发人员需要熟悉Android软件开发工具包(SDK)和本地开发工具包(NDK)。Android SDK和NDK分别参考https://developer.android.com/index.html和https://developer.android.com/ndk/index.html
在以后的章节中,所采用的开发平台默认是Android,开发人员有使用Android SDK和NDK的经验。Linux上的应用程序开发应该是类似的。
在Snapdragon平台上进行OpenCL开发有以下几个先决条件:
-
支持OpenCL的Snapdragon设备。并非所有的Snapdragon设备都支持OpenCL。详情请参阅表3-1。
-
OpenCL软件。OpenCL在Adreno gpu上依赖于QTI专有库。
- 检查设备是否安装了OpenCL库。
—核心库为libOpenCL.so,通常位于设备上的/vendor/lib目录下。
- 一些供应商可能选择不包含OpenCL软件(例如,谷歌的Nexus和Pixel设备)。
-
OpenCL必须运行在NDK层。
-
开发和测试不需要Root访问权限,但是在性能模式下运行soc可能需要Root访问权限。
4.2 调试工具
由于GPU执行的并行性,调试OpenCL内核通常很有挑战性。Adreno gpu支持内核中的printf函数,这对于调试非常有用。为了使用printf,建议减少工作负载,打印带有条件的变量,并避免打印太多的变量,因为printf会降低代码执行速度。例如,可以只启用有问题的工作组,甚至是单个有问题的工作项(通过在函数CLEnqueueNDRangeKernel中设置适当的偏移量)
了解设备的软件版本很重要,因为一些bug或问题可能已经在新版本中修复。要查询软件(驱动程序)和编译器版本,可以使用一个名为clGetDeviceInfo的API函数。
4.3 Snapdragon Profiler
Snapdragon Profiler是QTI提供的一个分析工具,运行在Windows, Mac和Linux平台上,允许开发者分析运行在Android上的Snapdragon处理器的CPU, GPU, DSP,memory,power, thermal, network等数据。它支持OpenCL和许多图形API,如OpenGL ES和Vulkan。详情请参:https://developer.qualcomm.com/software/snapdragon-profiler
下面是Snapdragon分析器为OpenCL分析提供的一些关键特性。
- profiler有一个内核分析器,允许开发人员对给定的内核进行静态分析。它提供了诸如寄存器占用、总指令数以及每种操作类型的指令数等信息,以帮助开发人员更好地优化内核。
- 分析器为给定的OpenCL应用程序提供OpenCL API跟踪和日志。它允许开发人员从API级别识别和解决瓶颈,并调试应用程序。
- 分析器提供诸如GPU占用率,ALU利用率,L1/L2缓存命中率等信息,这对于开发者识别内核的性能问题是至关重要的。
- 分析器支持基于命令行的应用程序,以及Android GUI应用程序。
4.4 性能分析
对于一个应用程序,准确地分析其性能是至关重要的。下面将讨论两种常用的方法,CPU定时器和GPU定时器,以及它们的主要区别。
4.1.1 CPU计时器
CPU计时器用于测量来自主机端的OpenCL调用的完整执行时间。这可以通过使用C/C++编程语言标准库中的任何日期和时间函数来实现。使用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函数可以分为阻塞调用和非阻塞调用。对于非阻塞调用,必须小心使用CPU定时器:
-
非阻塞调用是指主机在指令提交后继续执行下一条指令(通常在另一个CPU线程中排队等待执行),而不是等待函数调用完成。
- 内核执行API函数clEnqueueNDRangeKernel是一个非阻塞函数。
-
对于非阻塞调用,真正的执行时间不是函数调用之间的时间差。
当使用CPU计时器从主机端测量内核执行时间时,必须通过使用clWaitforEvent调用(如果非阻塞调用有一个事件ID)或clFinish来确保函数完成。同样的规则也适用于内存传输调用。
4.4.2 GPU计时器
所有OpenCL队列函数调用都可选择返回一个事件对象到主机,OpenCL分析API可以使用该对象来查询执行时间。Adreno GPU有自己的时钟和计时器来测量函数的执行流程,GPU的执行时间是由独立于操作系统的GPU硬件计数器决定的。
要启用GPU计时器功能,需要在当前命令队列的clCreateCommandQueue或clSetCommandQueueProperty的属性参数中设置CL_QUEUE_PROFILING_ENABLE标志。同样,一个事件对象必须提供给enqueue函数。函数完成后,API函数clGetEventProfilingInfo应使用将获取命令执行的概要信息。
对于一个clEnqueueNDRangeKernel调用,使用clGetEventProfilingInfo函数和四个配置参数,包括CL_PROFILING_COMMAND_(QUEUED, SUBMIT, START, END),可以在Adreno gpu中提供内核启动延迟和内核执行时间的准确图像,如图4-1。
- 前两个参数CL_PROFILING_COMMAND_(QUEUED和SUBMIT)之间的区别,给出了软件开销和CPU缓存操作开销的概念。OpenCL软件可能会选择先将内核队列化,然后再将其与队列中随后的几个内核一起提交,例如,当队列中的内核数量足够多时。开发人员可以使用clFlush函数来加速提交。
- CL_PROFILING_COMMAND_(提交和开始)之间的区别可以给出一个GPU正在处理的许多其他工作的想法。
- 内核在GPU上的实际执行时间是CL_PROFILING_COMMAND_(START和END)的差值。
开发人员应该专注于最小化实际的内核执行时间。与其他两个通常难以控制的计时器相比,这可以相对容易地改进。
4.4.3 GPU定时器和CPU定时器
由于GPU和CPU定时器都可以用来分析性能,哪一个应该用于应用程序?虽然GPU定时器可以精确测量GPU的执行时间,但是一些硬件操作(如缓存刷新)和一些软件操作(如CPU主机与GPU同步)不在GPU时钟系统中。因此,GPU定时器在内核执行中报告的性能数字可能比CPU定时器更好。以下是推荐的两种做法:
- GPU定时器应该用来测量内核优化。从GPU执行的角度来看,GPU计时器能够准确地告诉我们优化步骤能够实现多少优化。
- CPU定时器应该用于测量整个应用程序的端到端性能。如果OpenCL程序只是整个应用程序管道的一部分,这一点很重要。
4.4.4 性能模式
Snapdragon SOCs具有先进的动态时钟和电压控制机制,自动控制系统,使其在某些场景下运行在节电模式,以节省电池。通常情况下,如果工作负载较大,系统可能会自动提高时钟速率和电压,将设备推入所谓的性能模式,以提高性能,满足工作负载需求。
对于OpenCL优化,如果系统动态改变时钟速率,将很难理解和分析性能。因此,为了分析的一致性和准确性,建议启用性能模式。
在没有性能模式设置的情况下,序列中的第一个OpenCL内核通常显示更大的启动延迟和更慢的执行时间。在启动真正的GPU工作负载之前,可能需要使用一些简单的内核来预热GPU。
OpenCL内核的性能并不仅仅依赖于GPU。在CPU上运行的API函数和GPU上的内核执行一样重要。为了达到最佳性能,CPU和GPU都应该启用性能模式。另外,为了减少UI渲染的干扰,建议:
- 确保被分析的应用程序全屏呈现,以便没有其他活动正在更新屏幕。
- 如果它是一个本地应用程序,确保SurfaceFlinger不是在Android上运行。这可以确保CPU和GPU只被应用程序使用。
注意:对于Adreno A3x、A4x和A5x gpu,启用性能模式所需的命令序列略有不同。
4.4.5 GPU频率控制
应用程序可以利用cl_qcom_perf_hint扩展来控制GPU频率。这个扩展允许应用程序在创建OpenCL上下文时设置性能提示属性。性能级别分为“高”、“正常”和“低”。NORMAL性能级别允许动态时钟和电压控制。“HIGH”和“LOW”性能级别分别禁用动态控制,并强制GPU以其最大和最小频率运行。
注意:性能级别只是一个提示。驱动试图参照并遵从这些提示,但诸如thermal控制或外部应用程序或服务等因素可能会覆盖这些提示。性能提示扩展为应用程序在功率/性能权衡方面提供了一些灵活性。但是,应该谨慎使用,因为它对SOC级功耗有重大影响。