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-1 使用Adreno gpu进行OpenCL开发的关键需求

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-1 Adreno图形处理器中clEnqueueNDRange的调用配置标志

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级功耗有重大影响。