OpenCL 通用编程与优化（14）

最新推荐文章于 2023-12-13 13:48:20 发布

weixin_38498942

最新推荐文章于 2023-12-13 13:48:20 发布

阅读量920

点赞数

分类专栏： OpenCL sdk 文章标签： OpenCL

本文链接：https://blog.csdn.net/weixin_38498942/article/details/129180255

版权

sdk 同时被 2 个专栏收录

281 篇文章 30 订阅

订阅专栏

OpenCL

46 篇文章 22 订阅

订阅专栏

文章介绍了AdrenoGPU支持的OpenCL供应商扩展，包括性能提示、上下文创建优先级、可记录命令队列和受保护上下文等功能，旨在提高性能和效率，同时强调了使用这些扩展时需注意的可移植性和版本兼容性问题。

摘要由CSDN通过智能技术生成

9. Adreno GPU中的OpenCL扩展

OpenCL平台或设备可以通过扩展机制支持或公开未在核心标准中采用的功能，以进一步提高其可用性，公开硬件功能并创建差异化的体验。 OpenCL标准根据其采用状态支持几种类型的扩展名，并且供应商采用了它们：

KHR扩展名已由OpenCL标准工作组批准，但是供应商支持是可选的。代表这些扩展的字符串以cl_khr开头。
- 这些扩展程序通常具有多个供应商的支持，并且具有一致性测试供应商，如果他们声称支持，则必须通过。
- KHR 扩展的规范可在 Khronos 的官方 OpenCL 网站上找到。
- 一些 KHR 扩展可能成为更新版本的 OpenCL 标准中的核心功能。
EXT 扩展由 OpenCL 标准工作组批准，但对于供应商支持是可选的。表示这些扩展的字符串以 cl_ext 开头。
- 这些更多是出于实验目的，通常没有一致性测试。
- 这些扩展可能成为KHR甚至核心功能。
供应商扩展具有特定的语法，通常只在特定供应商的平台上运行。他们的名字通常包含供应商公司名称的首字母缩略词。
- 例如，Adreno GPU 中的 OpenCL 供应商扩展内部有“qcom”，例如 cl_qcom_accelerated_image_ops。

此外，还可能有私人或内部扩展，供应商可能不会公开，但可以向客户使用。扩展程序可以针对OpenCL平台或OpenCL设备：

可以通过cl_platform_info的参数设置为cl_platform_extensions的参数，可以通过API函数ClgetPlatFormInfo获得给定平台上可用的扩展。
OpenCL 设备可用的扩展可以通过 API 函数 clGetDeviceInfo 获取，设备信息参数设置为 CL_DEVICE_EXTENSIONS。

这两个功能都将返回可用扩展名的名称字符串列表。使用扩展时有警告：

扩展可以有不同的版本。
供应商扩展可以被供应商弃用或取消功能。
扩展可以成为 OpenCL 标准的核心功能。如果扩展不可用，请检查核心规范是否已采用它。
开发人员在使用扩展之前必须查询平台上可用的扩展列表，以获得更好的可移植性。

下表显示了Snapdragon 888设备支持的KHR和供应商扩展。扩展名的详细文档可在Adreno Opencl SDK中找到，可以从QTI开发人员网络网站（https://developer.qualcomm.com）下载。

在接下来的几节中，将跳过 KHR 扩展，并提供 Adreno GPU 中可用的供应商扩展的高级概述。

表 9-1 骁龙 888 设备支持的扩展列表

拓展类型	拓展名	参考
KHR extensions	cl_khr_3d_image_writes cl_img_egl_image cl_khr_byte_addressable_store cl_khr_depth_images cl_khr_egl_event cl_khr_egl_image cl_khr_fp16 cl_khr_gl_sharing cl_khr_global_int32_base_atomics cl_khr_global_int32_extended_atomics cl_khr_local_int32_base_atomics cl_khr_local_int32_extended_atomics cl_khr_image2d_from_buffer cl_khr_mipmap_image cl_khr_srgb_image_writes cl_khr_subgroups cl_khr_integer_dot_product cl_ext_image_from_buffer cl_ext_image_requirements_info	The OpenCL™ extension specification: https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html
Vendor extensions	cl_qcom_create_buffer_from_image cl_qcom_ext_host_ptr cl_qcom_ion_host_ptr cl_qcom_perf_hint cl_qcom_other_image cl_qcom_subgroup_shuffle cl_qcom_vector_image_ops cl_qcom_extract_image_plane cl_qcom_android_native_buffer_host_ptr cl_qcom_protected_context cl_qcom_priority_hint cl_qcom_compressed_image cl_qcom_ext_host_ptr_iocoherent cl_qcom_accelerated_image_ops cl_qcom_reqd_sub_group_size cl_qcom_recordable_queues cl_qcom_android_ahardwarebuffer_host_ptr cl_qcom_bitreverse cl_qcom_create_buffer_from_image cl_qcom_dot_product8 cl_qcom_ml_ops cl_qcom_extended_query_image_info cl_qcom_filter_bicubic cl_qcom_onchip_global_memory	Adreno Opencl SDK文档（检查Qualcomm开发人员网络）

9.1 OS依赖性供应商扩展

本节介绍扩展 OpenCL 运行时功能的扩展，无需 OpenCL 内核语言中的新功能。某些扩展可能取决于 Android 操作系统及其低级设计，这可能会发生变化。必须经常检查 Qualcomm 开发人员网络的 OpenCL SDK 以了解最新的扩展详细信息。本文档其他部分涵盖的扩展将被跳过。例如，多个零拷贝扩展的详细信息在第 7.4 节中。

9.1.1 性能提示 (cl_qcom_perf_hint)

此扩展允许应用程序在 OpenCL 上下文中请求设备所需的性能级别。更高的性能意味着设备上的更高频率。分别通过标志 CL_PERF_HINT_HIGH_QCOM、CL_PERF_HINT_NORMAL_QCOM 和 CL_PERF_HINT_LOW_QCOM 支持三个级别的性能提示，包括高、正常和低。

CL_PERF_HINT_HIGH_QCOM 是 OpenCL 上下文中设备的默认设置（从设备请求最高性能级别）。
CL_PERF_HINT_NORMAL_QCOM 是一个平衡的性能设置，由 GPU 频率和电源管理动态设置。
CL_PERF_HINT_LOW_QCOM 请求优先考虑低功耗的性能设置。

有两种方法可以使用此扩展：

使用 clCreateContext 函数创建上下文时，在上下文属性参数中指定性能提示标志。这是一个例子：

cl_context_properties properties[] =  {
    CL_CONTEXT_PERF_HINT_QCOM, CL_PERF_HINT_LOW_QCOM, 0}; clCreateContext(properties, 1, &device_id, NULL, NULL, NULL);

使用单独的 API 函数 clSetPerfHintQCOM 为现有上下文设置性能提示属性。此函数可用于设置或更新 CL_CONTEXT_PERF_HINT_QCOM 属性，而不管它是否在上下文时间指定为上下文属性之一。这是一个例子：

clSetPerfHintQCOM(context, CL_PERF_HINT_NORMAL_QCOM);

9.1.2 上下文创建的优先级提示 (cl_qcom_priority_hint)

此扩展允许应用程序指定所需的优先级，以便在 OpenCL 上下文中将排队的内核提交给设备。与性能提示标志一样，它定义了三个级别的优先级：

高优先级，CL_PRIORITY_HINT_HIGH_QCOM，这意味着排队的内核可以在其他具有较低优先级的上下文中排队之前提交给设备进行处理。
正常优先级，CL_PRIORITY_HINT_NORMAL_QCOM，默认行为。如果不使用此扩展，则选择将用于上下文的优先级。
低优先级，CL_PRIORITY_HINT_LOW_QCOM，与高优先级相反，内核在具有更高优先级的其他上下文中的其他内核之后提交给设备。

该提示应在上下文创建中使用clcreatecontext作为上下文属性提供。

9.1.3 可记录命令队列（cl_qcom_recordable_queues）

内核队列函数调用 clEnqueueNDRangeKernel 是将内核分派到 GPU 硬件的关键且要求很高的函数，因为它需要应用程序配置和验证许多内核参数，例如全局工作大小、工作组大小、事件依赖性等。对于内核必须排队进行重复处理的用例，例如视频处理应用程序，其中每一帧按顺序重复执行相同的内核，开发人员不能忽略不断设置参数的开销。

这个扩展没有重复这些操作，而是引入了一组新的 API 函数来记录内核排队的顺序。序列只需要生成一次，但可以分派多次。被记录序列中的任何内核的参数(称为记录)都可以在不重新记录整个命令序列的情况下被修改。因此，扩展可以节省CPU功率并提高调度延迟。这是此扩展的一些关键功能：

记录需要一个命令队列，必须使用CLCREATECOMMANDQUEUE使用可记录的属性CL_QUEUE_RECORDABLE_QCOM创建该命令队列。
必须使用称为ClNewRecordingQcom的函数创建记录对象。
要启动记录，请使用标准启用函数ClenqueuendrangeKernel，并使用ClNewRecordingQcom创建了此记录对象。
目前只能记录clEnqueueNDRangeKernel。
ClendRecordingQcom用于完成录制。
clEnqueueRecordingQCOM 用于将可记录队列中的所有内核排入队列以供 GPU 执行，这需要记录对象和“实时”命令队列。
- 该函数可用于更新可记录队列中内核的参数。
- 此功能的直播命令队列与记录的不同，必须是有序命令队列。
- 定义了一种机制来指定要更新的内核、要更改的参数列表以及要更改的每个参数的新值。
- 所有内核参数都可以更改，包括内核参数、全局大小等。
- 应用程序可以选择不更新记录。
此扩展不适用于内核排队内核 (KEK) 功能或 printf。
Khronos的OpenCL工作组一直在进行努力，以标准化一个KHR扩展名，该扩展名支持类似功能，同时更加通用和高级。
如果 OpenCL 工作组批准了 KHR 扩展，则可以弃用此供应商扩展。

9.1.4 cl_qcom_protected_context

此扩展允许应用程序创建所谓的受保护 OpenCL 上下文。在受保护的上下文中创建的 OpenCL 命令队列也被隐式地视为受保护的。受保护的 OpenCL 上下文支持使用特定 Qualcomm GPU 上可用的内容保护功能。此功能的主要目的是将内存分为受保护区域和不受保护区域，并防止将数据从受保护区域复制到不受保护区域。要使用此功能，必须创建具有 CL_CONTEXT_PROTECTED_QCOM 属性的上下文：

 cl_context_properties properties[] = {CL_CONTEXT_PROTECTED_QCOM, 1, 0};    protected_context = clCreateContext(properties, 1, &device_id, NULL, NULL, &err);

一旦创建了使用上下文的受保护命令队列，然后在整个应用程序中使用：

protected_queue = clCreateCommandQueue(protected_context, device_id, 0, &err);

基本上有两种方法可以在 Android 上创建受保护的内存对象：

受保护的图形缓冲区可以使用 GRALLOC_USAGE_PROTECTED 使用标志进行分配，并且可以在 OpenCL 中通过使用带有 clCreateBuffer 或 clCreateImage2D 的 cl_qcom_android_native_buffer_host_pointer 扩展来访问。
受保护的 ION 分配可以使用 ION_SECURE 标志从受保护的堆中创建，并且可以在 OpenCL 中通过使用带有 clCreateBuffer 或 clCreateImage2D 的 cl_qcom_ion_host_pointer 扩展来访问。

在这两种情况下，缓冲区都被OpenCL视为受保护的内存对象。一个具有一个或多个受保护的内存对象的OPENCL应用程序，因为参数只能在受保护的命令标题上进行。

cl_qcom_create_buffer_from_image

在 OpenCL 中，图像对象是一种不透明的数据结构，由 OpenCL API 中定义的函数进行管理。开发人员无法获得图像对象中数据存储方式的底层细节。与可以使用指针直接访问的缓冲区对象不同，必须在内核内部使用内置的图像读写函数（如 image_readf/image_writef）来访问图像数据。

在某些情况下，开发人员可能希望以 OpenCL C 语言中的原始指针的形式访问图像数据，而不是使用内置的图像读/写函数。此扩展可以使用例受益，如下所示：

通过OpenGL和OpenCL Interop扩展从间接地向OpenCL读取或写入EGL外部图像。
有经验的开发人员可能希望通过单个内存加载/存储操作读取/写入多个像素。
图像对象中的数据可能必须同时送入一个只接受它作为缓冲对象的内核，以及另一个只接受它作为图像对象的内核。

使用扩展程序，要从现有图像对象创建一个新的RAW BUFFER对象，显示了以下函数：
cl_mem clCreateBufferFromImageQCOM(cl_mem image, cl_mem_flags flags, cl_int *errcode_ret)

其中 image 是具有某些限制的有效图像。对图像类型、布局和并发读写访问有要求：

它支持除以下以外的所有图像：
- 图像类型 CL_MEM_OBJECT_IMAGE1D_BUFFER，或
- 使用 CL_MEM_USE_HOST_PTR 创建的图像。
数据布局：
- 返回的缓冲区引用为图像分配的数据存储并指向该数据存储中的原始像素。
- 数据布局等同于原点为 (0, 0, 0) 且区域为 (width, height, depth) 时 clEnqueueMapImage 生成的数据布局。
- 创建缓冲区的图像称为缓冲区的 parent_image。
要正确访问返回的缓冲区中的像素数据，客户必须使用clgetDeviceImageImageInfoqcom查询父映像的两个参数：
- 行距，使用CL_BUFFER_FROM_IMAGE_ROW_PITCH_QCOM。
- 切片间距，使用 CL_BUFFER_FROM_IMAGE_SLICE_PITCH_QCOM。
读写并发：
- 未定义：
  - 并发读取和写入缓冲区对象及其 parent_image 是未定义的。
  - 从同一parent_image创建的同时读取并写入对缓冲对象的读写不确定。
- 受支持的：
  - 仅从缓冲对象及其parent_image对象同时读取。
  - 定义了从同一图像创建的多个缓冲区对象的并发读取。
Khronos的OpenCL工作组一直在进行努力，以标准化一个KHR扩展名，该扩展名支持类似功能，同时更加通用和高级。
- 如果 OpenCL 工作组批准了 KHR 扩展，则可以弃用此供应商扩展。

有关如何使用它的更多详细信息和示例，请参阅Adreno Opencl SDK中的扩展文档。

9.1.6 cl_qcom_onchip_global_memory

此扩展提供了创建 OpenCL 缓冲区和图像的功能，这些缓冲区和图像驻留在 Adreno GPU 上更快访问的片上内存（以下称为 onchip_global_memory）中。一旦创建，这些对象就可以在 OpenCL 内核中使用，就像使用常规全局内存对象一样。

如果 cl_qcom_other_image 存在，此扩展进一步允许从在 onchip_global_memory 中创建的 OpenCL 缓冲区创建平面图像。

默认情况下，onchip_global_memory的内容在内核退出后不保留。然而，应用程序可以使用onchip_global_memory通过利用cl_qcom_recordable_queues扩展在两个或多个内核之间传递数据。onchip_global_memory的内容将在记录的整个队列期间有效。因此，应用程序可以使用onchip_global_memory将一个内核的输出链接到下一个内核的输入，只要这些内核是同一记录的一部分。Onchip_global_memory不会跨相同记录的不同队列保存。记录中的第一个内核必须将onchip_global_memory视为未初始化的。

通过使用在 onchip_global_memory 中分配的缓冲区和图像代替全局分配，应用程序可以实现功率和性能改进。一个例子是将 onchip_global_memory 用于内核管道中的中间缓冲区。

有关扩展文档和演示 onchip_global_memory 用法的示例，请参阅 Adreno OpenCL SDK。