GPU 优化技术-OpenCL 运行时 API 介绍
版权声明 本文首发于知乎专栏：移动端算法优化本专栏「移动端算法优化」所有文章著作权归作者所有。
允许个人直接分享本专栏文章到个人微博、朋友圈。但媒体（包括但不限于网站、微信公众号、微博营销号）转载需事先征得专栏同意。转载需在正文开头显著位置注明出处，给出原始链接，注明「发表于知乎专栏《移动端算法优化》」，并不得进行任何形式的修改演绎。<hr/>一、概述
OpenCL 作为一套通用异构平台编程框架，由两个部分组成：其一是在主机处理器（Host）执行的运行时 API；其二是基于 C99 标准扩展的 OpenCL C 语言，用于编写在设备处理器（OpenCL device）运行的内核（kernel）代码。
Host 端的运行时 API 负责管理资源，控制 host 和 device 端程序执行，构成了 OpenCL 程序的框架。不恰当地使用 OpenCL 运行时 API 可能带来极大的性能损失，甚至造成程序崩溃。
接下来我们会对 OpenCL 运行时 API 做系统而详细的介绍。首先给出 OpenCL 运行时 API 的整体编程流程图，然后针对每一步结合实例说明，最后展示基于 OpenCL 的图像转置代码并计算加速比。
二、OpenCL 运行时 API 的编程流程
2.1 OpenCL Host 端编程全流程
下图展示了一张典型的 OpenCL Host 端编程流程图。接下来我们会解释每一步的作用，列出相关的API并展示实例代码。
基于 OpenCL 运行时 API 的 Host 编程流程
使用 OpenCL API 编程务必及时检查 API 的返回值是否为 CL_SUCCESS。为了简化篇幅，第二节示例代码不检查 API 返回值，省略 malloc 内存释放。
2.2 OpenCL Host端编程分步详解
2.2.1 获取平台Platform
作用说明
Host 端编程第 1 步是获取硬件平台 Platform，查询 OpenCL 版本等平台信息。使用 clGetPlatformIDs 和 clGetPlatformInfo 两个API。
API函数说明
[*]clGetPlatformIDs
cl_int clGetPlatformIDs( cl_uint num_entries,
cl_platform_id *platforms,
cl_uint *num_platforms)
功能描述：
OpenCL 使用 cl_platform_id 表示平台，通过 clGetPlatformIDs 获取可使用的平台数组。当平台数未知，第一次调用 clGetPlatformIDs 获取平台个数，第二次调用 clGetPlatformIDs 获取平台对象。
参数说明：
- num_entries ：要获取的平台数量，如果 platforms 非空，则 num_entries 不能为 0。
- platforms ：返回获取的平台对象数组。
-num_platforms ：用于查询返回可用的平台数目，num_platforms 可设为 NULL 忽略。
- Return ：正常执行返回 CL_SUCCESS，异常返回值请参考中 4.1节。
[*]clGetPlatformInfo
cl_int clGetPlatformInfo( cl_platform_id platform,
cl_platform_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
功能描述：
获取平台相关信息，如 OpenCL 配置版本。
参数说明：
- platform ：查询的平台对象。
- param_name ：表示平台查询参数的枚举常量，参考表 2-1。
- param_value_size ：param_value 指向内存空间的字节数。
- param_value ：指向返回查询参数结果的内存指针。
- param_value_size_ret ：返回查询参数的实际长度。
- Return ：正常执行返回 CL_SUCCESS，异常返回值请参考中 4.1 节。
表2-1 OpenCL平台属性表
实例代码
下面给出两步查询法获取平台及相关信息的实例代码。
void PrintPlatformMsg(cl_platform_id *platform, cl_platform_info platform_info,
const char *platform_msg)
{
size_t size;
int err_num;
// 1. 第一步通过size获取打印字符串长度
err_num = clGetPlatformInfo(*platform, platform_info, 0, NULL, &size);
char *result_string = (char *)malloc(size);
// 2. 第二步获取平台信息到result_string
err_num = clGetPlatformInfo(*platform, platform_info, size, result_string, NULL);
printf(&#34;%s=%s\n&#34;, platform_msg, result_string);
free(result_string);
result_string = NULL;
}
cl_int err_num;
cl_uint num_platform;
cl_platform_id *platform_list;
// 1. 第一次调用获取平台数
err_num = clGetPlatformIDs(0, NULL, &num_platform);
printf(&#34;num_platform=%d\n&#34;, num_platform);
platform_list = (cl_platform_id *)malloc(sizeof(cl_platform_id) * num_platform);
// 2. 第二次调用获取平台对象数组
err_num = clGetPlatformIDs(num_platform, platform_list, NULL);
printf(&#34;err_num = %d\n&#34;, err_num);
// 打印平台信息
PrintPlatformMsg(&platform_list, CL_PLATFORM_PROFILE, &#34;Platform Profile&#34;);
PrintPlatformMsg(&platform_list, CL_PLATFORM_VERSION, &#34;Platform Version&#34;);
PrintPlatformMsg(&platform_list, CL_PLATFORM_NAME, &#34;Platform Name&#34;);
PrintPlatformMsg(&platform_list, CL_PLATFORM_VENDOR, &#34;Platform Vendor&#34;);
在高通865平台运行结果如下：
num_platform=1
err_num = 0
Platform Profile=FULL_PROFILE
Platform Version=OpenCL 2.0 QUALCOMM build: commit #d970ca5f2e changeid #Ifead41f47e Date: 07/14/21 Wed Local Branch:Remote Branch:
Platform Name=QUALCOMM Snapdragon(TM)
Platform Vendor=QUALCOMM2.2.2 获取设备Device
作用说明
获取平台下的 OpenCL 设备 Device，查询设备的硬件参数。使用 clGetDeviceIDs 和 clGetDeviceInfo两个API。
API函数说明
[*]clGetDeviceIDs
cl_int clGetDeviceIDs (cl_platform_id platform,
cl_device_type device_type,
cl_uint num_entries,
cl_device_id *devices,
cl_uint *num_devices)
功能描述：
获取平台可使用的 Device 对象数组。OpenCL 使用 cl_device_id 表示 Device 对象。
参数说明：
- platform ：clGetPlatformIDs 获取的 Platform ID。
- device_type ：获取 OpenCL Device 的类型，参考表 2-2。
- num_entries ：要获取的设备数量。
- devices ：返回获取的设备对象数组。
- num_devices：返回平台连接 device_type 类型设备数目，可设为 NULL 忽略。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 4.2 节。
表2-2 OpenCL设备类型列表
[*]clGetDeviceInfo
cl_int clGetDeviceinfo(cl_device_id device,
cl_device_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
功能描述：
获取设备相关信息，例如并行计算单元数，全局内存大小等等。
参数说明：
- device ：clGetDeviceIDs 获取的 Device ID。
- param_name ：表示设备查询参数的枚举常量。
- param_value_size ：param_value 指向内存空间的字节数。
- param_value ：指向返回查询参数结果的内存指针。
- param_value_size_ret：返回查询参数的实际长度。
- Return ：正常执行返回 CL_SUCCESS，异常返回值请参考中 4.2 节。
clGetDeviceInfo 与 clGetPlatformInfo的使用方式基本一致，由于设备属性较多，表 2-3 仅列出部分常用设备硬件属性。设备属性全表请参考中表 4.3。
表 2-3 OpenCL 部分设备属性列表
实例代码
下面给出获取设备及设备参数的实例代码。
cl_uint num_device;
cl_device_id device;
// 1. 获取平台GPU类型OpenCL设备的数量
err_num = clGetDeviceIDs(platform_list, CL_DEVICE_TYPE_GPU, 0, NULL, &num_device);
printf(&#34;GPU num_device=%d\n&#34;, num_device);
// 2. 获取一个GPU类型的OpenCL设备
err_num = clGetDeviceIDs(platform_list, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
// 对于cl_uint cl_ulong等返回类型参数只需要一步查询
cl_uint max_compute_units;
// 获取并打印OpenCL设备的并行计算单元数量
err_num = clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint),
&max_compute_units, NULL);
printf(&#34;max_compute_units=%d\n&#34;, max_compute_units);
cl_ulong global_mem_size;
// 获取并打印OpenCL设备的全局内存大小
err_num = clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong),
&global_mem_size, NULL);
printf(&#34;global_mem_size=%ld\n&#34;, global_mem_size);
size_t *p_max_work_item_sizes=NULL;
size_t size;
// CL_DEVICE_MAX_WORK_ITEM_SIZES表示work_group每个维度的最大工作项数目
// 1. 返回类型是size_t[]，首先查询返回信息的大小
err_num = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, 0, NULL, &size);
p_max_work_item_sizes = (size_t *)malloc(size);
// 2. 申请空间后查询结果并打印
err_num = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, size, p_max_work_item_sizes, NULL);
for (size_t i = 0; i < size / sizeof(size_t);i++)
{
printf(&#34;max_work_item_size_of_work_group_dim %zu=%zu\n&#34;, i, p_max_work_item_sizes);
}
在高通865平台运行结果如下
GPU num_device=1
max_compute_units=3
global_mem_size=3988809728
max_work_item_size_of_work_group_dim 0=1024
max_work_item_size_of_work_group_dim 1=1024
max_work_item_size_of_work_group_dim 2=10242.2.3 创建上下文Context
作用说明
选择获取的设备，创建上下文 Context。使用的API是 clCreateContext。OpenCL 使用 Context 管理命令队列、程序内核、内存等资源对象。
API函数说明
[*]clCreateContext
cl_context clCreateContext(const cl_context_properties *properties,
cl_uint num_devices,
const cl_device_id *devices,
void ( CL_CALLBACK *pfn_notify)(const char *errinfo, const void *private_info, size_t cb, void *user_data),
void *user_data,
cl_int *errcode_ret)
功能描述：
基于获取的平台和设备，创建一个 OpenCL 上下文，类型是 cl_context。
参数说明：
- properties ：上下文属性数组，每一项属性包含一个枚举常量，一个属性值，数组以 0 结尾。properties 指定了创建 Context 基于的 Platform，也可以设为 NULL，程序实现时自行选择 Platform。
- num_devices：devices 中指定的设备数。
- devices ：clGetDeviceIDs 返回的设备对象数组。
- pfn_notify ：注册回调函数，当 OpenCL 创建上下文失败时会执行回调函数。没有回调可设为 NULL
- user_data：传递给回调函数 pfn_notify 的指针参数，可设为 NULL。
- errcode_ret ：返回错误码，如果 errcode_ret 设为 NULL 不再返回错误码。
- Return ：OpenCL 上下文成功创建时，返回创建的 cl_contex t对象，errcode_ret 返回CL_SUCCESS。创建失败时返回 NULL，errorcode_ret 返回错误码。实例代码
下面的代码使用高通 865 平台的 GPU 设备创建 Context。
cl_context_properties context_prop = {0};
context_prop = CL_CONTEXT_PLATFORM;
context_prop = (cl_context_properties)platform_list;
context = clCreateContext(context_prop, 1, &device, NULL, NULL, &err_num);
if (err_num != CL_SUCCESS)
{
printf(&#34;Create Context failed with code=%d!\n&#34;, err_num);
}
else
{
printf(&#34;Context successfully created!\n&#34;);
}
在高通 865 平台运行结果如下：
Context successfully created!2.2.4 创建命令队列 CommandQueue
作用说明
为单个设备创建命令队列，使用的 API 是 clCreateCommandQueueWithProperties。操作命令入队后依据队列属性顺序或者乱序执行。
API函数说明
[*]clCreateCommandQueueWithProperties
cl_command_queue
clCreateCommandQueueWithProperties(cl_context context,
cl_device_id device,
const cl_queue_properties *properties,
cl_int *errcode_ret)
功能描述：
基于 Context 和唯一的 Device，按照 properties 属性创建命令队列 cl_command_queue。
参数说明：
- context ：有效的上下文对象 cl_context。
- device ：与 context 关联的设备。
- properties ：命令队列属性数组，每一项属性包含一个枚举常量，一个属性值，数组以 0 结尾。参考表 2-4。
- errcode_ret：返回错误码，如果 errcode_ret 设为 NULL 不再返回错误码。
- Return ：OpenCL 命令队列成功创建时，返回创建的 cl_command_queue 对象，errcode_ret 返回 CL_SUCCESS。创建失败时返回NULL，errorcode_ret 返回错误码。
表 2-4 命令队列属性表
实例代码
下面的代码创建一个 Host 端入队的命令队列，设置 CL_QUEUE_PROFILING_ENABLE 属性用于性能分析。
// OpenCL设备命令执行分为入队、提交、启动、结束和完成5个时间点，创建命令队列时使能CL_QUEUE_PROFILING_ENABLE才能获取设备记录的相应时间。
cl_command_queue_properties queue_prop[] = {CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE, 0};
cl_command_queue command_queue = NULL;
command_queue = clCreateCommandQueueWithProperties(context, device, queue_prop, &err_num);
if (err_num != CL_SUCCESS)
{
printf(&#34;Create CommandQueue failed with code=%d!\n&#34;, err_num);
}
else
{
printf(&#34;Host in-order profiling CommandQueue successfully created!\n&#34;);
}
在高通 865 平台运行结果如下
Host in-order profiling CommandQueue successfully created!2.2.5 创建并编译内核程序 Program
作用说明
这一步通过 OpenCL C 源码字符串或程序二进制两种方式之一创建内核程序 Program。编译 Program 生成二进制，检查编译错误并获取二进制代码。使用二进制代码创建 Program 能显著减少编译时间。
API函数说明
[*]clCreateProgramWithSource
cl_program clCreateProgramWithSource(cl_context context,
cl_uint count,
const char **strings,
const size_t *lengths,
cl_int *errcode_ret)
功能描述：
使用 OpenCL C 源代码创建 cl_program，程序对象关联 context 中的所有设备。
参数说明：
- context ：有效的Context对象。
- count ：表示 strings 中字符串的个数。
- strings ：字符串数组指针，所有的字符串构成设备源代码。
- lengths ：表示 strings 每个字符串的长度。lengths 可以设为 NULL，字符串以 0 结尾自动计算长度。
- errcode_ret：返回错误码。errcode_ret 设为 NULL 则不再返回错误码。
- Return ：Program 对象成功创建时，返回创建的 cl_program 对象，errcode_ret 返回CL_SUCCESS。创建失败时返回 NULL，errorcode_ret 返回错误码。
[*]clCreateProgramWithBinary
cl_program clCreateProgramWithBinary(cl_context context,
cl_uint num_devices,
const cl_device_id *device_list,
const size_t *lengths,
const unsigned char **binaries,
cl_int *binary_status,
cl_int *errcode_ret)
功能描述：
向设备载入构建好的可执行程序二进制并创建 cl_program。
参数说明：
- context : 有效的 Context 对象。
- num_devices :device_list 中设备个数。
- device_list ：context 关联的设备数组。程序二进制需要载入 device_list 所列出的设备中，因此不能为NULL。
- lengths ：binaries 数组中每个二进制文件的长度。
- binaries ：二进制文件数组。对于设备device_list，其程序二进制文件是 binaries，文件的长度是 lengths，三者一一对应。
- binary_status：返回每个设备对应的二进制是否成功加载。成功加载返回 CL_SUCCESS。binary_status可设为NULL以忽略。
- errcode_ret ：返回错误码。errcode_ret 设为NULL则不再返回错误码。
- Return ：Program 对象成功创建时，返回创建的 cl_program 对象，errcode_ret 返回CL_SUCCESS。创建失败时返回NULL，errorcode_ret 返回错误码。
[*]clBuildProgram
cl_int clBuildProgram(cl_program program,
cl_uint num_devices,
const cl_device_id *device_list,
const char *options,
void (CL_CALLBACK *pfn_notify)( cl_program program, void *user_data),
void *user_data)
功能描述：
使用 Program 源代码或者二进制生成设备可执行程序。
参数说明：
- program ：创建的程序对象。
- num_devices ：device_list 中的设备数目。
- device_list ：program 关联的设备对象数组。device_list 设为 NULL 时，为 program 关联的所有设备编译可执行程序。device_list 非空则仅为 device_list 中给出的设备编译可执行程序。
- options ：表示编译选项的字符串，例如按 OpenCL2.0 标准编译，-cl-std=CL2.0。请参考中5.8.4节 CompilerOptions。
- pfn_notify ：注册回调函数。如果 pfn_notify 非空，clBuildProgram 在编译开始后立即返回，生成结束时异步调用回调函数。如果 pfn_notify 为空，clBuildProgram 必须等待生成结束才能返回。
- user_data：传递给回调函数的参数，可以设置为 NULL。
- Return ：成功生成可执行程序返回 CL_SUCCESS，错误码参考中 5.8.2 节。
clBuildProgram 包含编译和链接过程，OpenCL2.0 也可以单独调用 clCompileProgram 编译，调用clLinkProgram 链接。
[*]clGetProgramBuildInfo
cl_int clGetProgramBuildInfo(cl_program program,
cl_device_id device,
cl_program_build_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
功能描述：
查询 Program 对象关联设备的编译信息，尤其是编译错误信息。
参数说明：
- program ：查询的程序对象。
- device ：指定查询编译信息的设备，设备必须与program 关联。
- param_name ：表示编译信息查询参数的枚举常量，参考表 2-5。
- param_value_size ：param_value 指向内存空间的字节数。
- param_value ：指向返回查询参数结果的内存指针。
- param_value_size_ret：返回查询参数的实际长度。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中 5.8.7 节。
表 2-5 Program 编译信息查询表
[*]clGetProgramInfo
cl_int clGetProgramInfo(cl_program program,
cl_program_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
功能描述：
查询程序信息，例如获取 clCreateProgramWithBinary 所需的可执行二进制。
参数说明：
- program ：查询的程序对象。
- param_name ：表示编译信息查询参数的枚举常量，参考表 2-6。
- param_value_size ：param_value 指向内存空间的字节数。
- param_value ：指向返回查询参数结果的内存指针。
- param_value_size_ret：返回查询参数的实际长度。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中 5.8.7 节。
表 2-6 程序属性查询表
实例代码
下面代码使用 OpenCL C 代码创建 Program，编译生成后将可执行二进制保存到 bin 文件。
// 1. 读取OpenCL C源代码
char *source = ClUtilReadFileToString(prog_name);
cl_int err = 0;
// 2. 使用源代码创建program
program = clCreateProgramWithSource(context, 1, (const char **)&source, NULL, &err);
// 3. 使用OpenCL2.0标准编译Program
err |= clBuildProgram(program, 0, NULL, &#34;-cl-std=CL2.0&#34;, NULL, NULL);
if (err != CL_SUCCESS)
{
// 如果编译失败，获取并打印错误信息
fprintf(stderr, &#34;Error %d with clBuildProgram.&#34;, err);
static const size_t LOG_SIZE = 2048;
char log;
log = 0;
err = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, LOG_SIZE, log, NULL);
if (CL_INVALID_VALUE == err)
{
fprintf(stderr, &#34;There was a build error, but there is insufficient space allocated to &#34;
&#34;show the build logs.\n&#34;);
}
else
{
fprintf(stderr, &#34;Build error:\n%s\n&#34;, log);
}
exit(-1);
}
else
{
// 4. 打印编译成功信息
printf(&#34;Program built Ok!\n&#34;);
}
cl_uint num_devices;
// 5. 获取程序关联设备数
err |= clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &num_devices, NULL);
// 6. 获取程序关联设备ID
cl_device_id *p_devices = (cl_device_id *)malloc(sizeof(cl_device_id) * num_devices);
err |= clGetProgramInfo(program, CL_PROGRAM_DEVICES, sizeof(cl_device_id) * num_devices,
p_devices, NULL);
// 7. 获取设备程序二进制代码长度
size_t *p_program_binary_sizes = (size_t *)malloc(sizeof(size_t) * num_devices);
// 8. 获取设备程序二进制代码
err |= clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t) * num_devices,
p_program_binary_sizes, NULL);
cl_uchar **p_program_binaries = (cl_uchar **)malloc(sizeof(cl_uchar *) * num_devices);
for (cl_uint i = 0; i < num_devices; i++)
{
p_program_binaries = (cl_uchar *)malloc(p_program_binary_sizes);
printf(&#34;Binary size for device %d=%zu\n&#34;, i, p_program_binary_sizes);
}
err |= clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(cl_uchar *) * num_devices,
p_program_binaries, NULL);
if (err != CL_SUCCESS)
{
printf(&#34;Error Occur!\n&#34;);
}
// 9. 保存可执行二进制代码到文件
for (cl_uint i = 0; i < num_devices; i++)
{
char fname;
sprintf(fname, &#34;Device%dProg.bin&#34;, i);
ClUtilWriteStringToFile(p_program_binaries, p_program_binary_sizes, fname);
printf(&#34;Wrote file %s\n&#34;, fname);
}
在高通865平台运行结果如下
Program built Ok!
Binary size for device 0=3516
Wrote file Device0Prog.bin2.2.6 创建内核对象Kernel
作用说明
第 6 步使用生成好的 Program 对象创建内核对象 kernel，类型为 cl_kernel。查询 kernel 的工作组属性以设置第 9 步执行参数。
API函数说明
[*]clCreateKernel
cl_kernel clCreateKernel(cl_program program,
const char *kernel_name,
cl_int *errcode_ret)
功能描述：
根据内核函数名，从 program 对象创建 kernel 对象。
参数说明：
- program ：已经生成可执行二进制的内核程序对象。
- kernel_name ：program 中以 __kernel 修饰的函数名。
- errcode_ret：返回错误码，如果 errcode_ret 设为 NULL 不再返回错误码。
- Return ：Kernel对象成功创建时，返回创建的 cl_kernel 对象，errcode_ret 返回CL_SUCCESS。创建失败时返回 NULL，errorcode_ret 返回错误码。
[*]clGetKernelWorkGroupInfo
cl_int clGetKernelWorkGroupInfo(cl_kernel kernel,
cl_device_id device,
cl_kernel_work_group_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
功能描述：
返回指定设备kernel 对象的工作组信息，例如最大工作组尺寸。
参数说明：
- kernel ：查询的内核对象。
- device ：与 kernel 关联的指定 Device ID。kernel 只与单个设备关联时，可以设为 NULL。
- param_name ：表示工作组信息查询参数的枚举常量，参考表 2-7。
- param_value_size ：param_value 指向内存空间的字节数。
- param_value ：指向返回查询参数结果的内存指针。
- param_value_size_ret：返回查询参数的实际长度。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中 5.9.3 节。
表 2-7 Kernel 工作组属性查询表
实例代码
cl_kernel kernel;
// 1. 创建Kernel
kernel = clCreateKernel(program, kernel_name, &err_num);
if (err_num != CL_SUCCESS)
{
printf(&#34;create kernel failed.\n &#34;);
return NULL;
}
size_t max_work_group_size;
size_t perferred_work_group_size_multiple;
err_num = clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t),
&max_work_group_size, NULL);
err_num |=
clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(size_t), &perferred_work_group_size_multiple, NULL);
if (err_num != CL_SUCCESS)
{
printf(&#34;Get kernel info failed.\n &#34;);
return NULL;
}
printf(&#34;Kernel %s max workgroup size=%zu\n&#34;, kernel_name, max_work_group_size);
printf(&#34;Kernel %s perferred workgroup size multiple=%zu\n&#34;, kernel_name,
perferred_work_group_size_multiple);
基于第三节转置 kernel，在高通 865 平台运行结果如下：
Kernel TransposeKernel max workgroup size=1024
Kernel TransposeKernel perferred workgroup size multiple=1282.2.7 创建内存对象
作用说明
创建内存对象并读入数据。OpenCL Buffer 对象用于存储一维的标量、向量或自定义结构体数据。Image 对象则专门用于保存一到三维的图像数据。OpenCL 通过这两种内存对象实现 Host 和 Device 之间大量数据交换。
2.2.7.1 OpenCL Buffer API 说明
本节介绍 Buffer 对象的创建和读写，给出实例代码。
[*]clCreateBuffer
cl_mem clCreateBuffer(cl_context context,
cl_mem_flags flags,
size_t size,
void *host_ptr,
cl_int *errcode_ret)
功能描述：
创建 Buffer 对象。
参数说明：
- context ：为context 对象分配 buffer 对象。
- flags ：以组合 bit 位枚举常量的方式，指定 buffer 的分配和使用信息，参考表 2-8。
- size ：申请内存空间字节数。
- host_ptr ：应用程序在 Host 端已经申请的内存空间指针。
- errcode_ret：返回错误码。errcode_ret 设为 NULL 则不再返回错误码。
- Return ：buffer 对象成功创建时，返回创建的 cl_mem 对象，errcode_ret 返回 CL_SUCCESS。创建失败时返回 NULL，errorcode_ret 返回错误码。
表 2-8 OpenCL Mem 对象读写 flag 列表
[*]clEnqueueRead/WriteBuffer
cl_int clEnqueueReadBuffer (cl_command_queue command_queue,
cl_mem buffer,
cl_bool blocking_read,
size_t offset,
size_t size,
void *ptr,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event)
cl_int clEnqueueWriteBuffer(cl_command_queue command_queue,
cl_mem buffer,
cl_bool blocking_write,
size_t offset,
size_t size,
const void *ptr,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event)
功能描述：
clEnqueueReadBuffer 从 buffer 对象读取数据到 Host 端内存，clEnqueueWriteBuffer 将 Host 内存数据写入 buffer 对象。
参数说明：
- command_queue ：读写指令入队的命令队列对象。命令队列和 buffer 应在同一 context 下创建。
- buffer ：有效的 buffer 对象。
- blocking_write和blocking_read：如果设置为 CL_TRUE，表示阻塞式操作，数据传输结束才能返回。CL_FALSE表示非阻塞操作，命令入队立即返回。
- offset ：读写 buffer object 的字节偏移量。
- size ：读写数据的字节数。
- ptr ：host 端内存缓冲区地址。
- num_events_in_wait_list ：event_wait_list 内事件的数目。
- event_wait_list ：cl_event 数组。执行内核操作前，需要等待 event_wait_list 内事件执行完成。如果不需要等待，event_wait_list 设为 NULL，num_events_in_wait_list 设为 0。
- event ：如果非 NULL，则此参数用于查询命令执行状态，或等待命令完成。clEnqueue 指令普遍包含 num_events_in_wait_list，event_wait_list，event 这三个参数。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中 5.2.2 节。Buffer 操作实例代码
int num_of_elements = 512;
int err = 0;
cl_uchar *h_A = (cl_uchar *)malloc(num_of_elements);
cl_uchar *h_B = (cl_uchar *)malloc(num_of_elements);
for (int i = 0; i < num_of_elements; i++)
{
h_A = i % 0xFF;
}
cl_mem buffer_A =
clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_uchar) * num_of_elements, NULL, &err);
cl_mem buffer_B =
clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_uchar) * num_of_elements, NULL, &err);
// write from host ptr to buffer
err = clEnqueueWriteBuffer(cmd_queue, buffer_A, CL_TRUE, 0, sizeof(cl_uchar) * num_of_elements, h_A,
0, NULL, NULL);
/* Run a device kernel copy data from buffer_A to buffer_B */
// Read from buffer to host ptr
err = clEnqueueReadBuffer(cmd_queue, buffer_B, CL_TRUE, 0, sizeof(cl_uchar) * num_of_elements, h_B, 0,
NULL, NULL);
clReleaseMemObject(buffer_A);
clReleaseMemObject(buffer_B);
// Return true if two arrays match
bool match = CompareArray(h_A, h_B, num_of_elements);
if (match)
printf(&#34;A and B match!\n&#34;);
else
printf(&#34;A and B mismatch!\n&#34;);
free(h_A);
free(h_B);
高通865平台运行结果如下
A and B match!2.2.7.2 OpenCL Image API 说明
Image 对象封装了图像大小、图像格式、坐标模式、插值模式等多种信息。Image 对象在 Device 端可使用采样器 Sampler 方便地读取图像。本节介绍 Image 对象和采样器对象的创建，Image 对象的读写并给出实例代码。
[*]clCreateImage
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)
功能描述：
创建指定格式、类型、尺寸的图像。
参数说明：
- context ：为 context 对象分配 Image 对象。
- flags ：指定缓冲区的分配和使用信息，枚举类型与 clCreateBuffer 一致。
- image_format：图像格式，包括图像通道顺序和数据类型。
- image_desc ：图像描述，包括图像类型、图像宽高和 pitch 等参数。
- host_ptr ：host端内存地址，可用于初始化图像数据或设为 NULL。
- errcode_ret ：返回错误码。errcode_ret 设为 NULL 则不再返回错误码。
- Return ：buffer 对象成功创建时，返回创建的 cl_mem 对象，errcode_ret 返回 CL_SUCCESS。创建失败时返回 NULL，errorcode_ret 返回错误码。
cl_image_format 结构体包含image_channel_order 和image_channel_data_type 两个成员，详情查询中表 5.6 和 5.7。举例来说，image_channel_order=CL_RGBA，image_channel_data_type=CL_UNSIGNED_INT8表示图像每个像素包括 RGBA 四个通道，每个通道的数据都是 8 位无符号整数。内存布局如下。
cl_image_desc 结构体包含图像类型，图像尺寸和一个 cl_mem 对象，参考中 5.3.1.2 节。
以 2D 图像为例，关键的尺寸参数 image_height，image_width 和 image_row_pitch。其中 image_height 和 image_width 以像素为单位。为了内存对齐，往往在图像每一行结尾填充位数，image_row_pitch 表示图像每一行的字节数。
// 图像format 根据具体的数据确定
cl_image_format image_format;
image_format.image_channel_order = CL_RGBA;
image_format.image_channel_data_type = CL_UNSIGNED_INT8;
cl_image_desc image_desc;
cl_uint pitch_align;
// 查询设备的2Dimage对齐字节
err_num = clGetDeviceInfo(device, CL_DEVICE_IMAGE_PITCH_ALIGNMENT, sizeof(cl_uint),
&pitch_align, NULL);
printf(&#34;image2D pitch align=%d\n&#34;, pitch_align);
memset(&image_desc, 0, sizeof(cl_image_desc));
image_desc.image_type = CL_MEM_OBJECT_IMAGE2D;
image_desc.image_height = height;
image_desc.image_width= width;
// 图像每一行的字节数，宽度*通道*元素大小，做内存对齐
image_desc.image_row_pitch = (width*channels*sizeof(cl_uchar) + pitch_align - 1) / pitch_align * pitch_align;
高通865平台的查询结果是
image2D pitch align=64
[*]clEnqueueRead/WriteImage
cl_int clEnqueueReadImage(cl_command_queue command_queue,
cl_mem image,
cl_bool blocking_read,
const size_t *origin,
const size_t *region,
size_t row_pitch,
size_t slice_pitch,
void *ptr,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event)
cl_int clEnqueueWriteImage(cl_command_queue command_queue,
cl_mem image,
cl_bool blocking_write,
const size_t *origin,
const size_t *region,
size_t input_row_pitch,
size_t input_slice_pitch,
const void * ptr,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event)
功能描述：
clEnqueueReadImage 从 image 对象读取数据到 host 内存，clEnqueueWriteImage 将 host 内存数据写入 image 对象。
参数说明：
- command_queue ：读写指令入队的命令队列对象。命令队列和 image 应在同一 context 下创建
- image ：有效的 image 或 image array 对象。
- blocking_read/blocking_write：如果设置为 CL_TRUE，表示阻塞式操作，数据传输结束才能返回。 CL_FALSE 表示非阻塞操作，命令入队立即返回。
- origin ：定义 image 对象的三维原点坐标 (x,y,z)，用于控制像素数据读写偏移。
- region ：定义 image 对象的 (width,height,depth) 坐标范围。因此图像数据读写的坐标范围是 (x,y,z) 到 (x+width,y+height,z+depth)。
- row_pitch ：表示 2D 以上图像每行的字节数。
- slice_pitch ：表示 3D 图像每层的字节数。
- ptr ：host 端内存缓冲区地址。
- num_events_in_wait_list ：event_wait_list 内事件的数目。
- event_wait_list ：cl_event 数组。执行内核操作前，需要等待 event_wait_list 内事件执行完成。如果不需要等待，event_wait_list 设为NULL，num_events_in_wait_list 设为 0。
- event ：如果非 NULL，则此参数用于查询命令执行状态，或等待命令完成。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中 5.3.3 节。
[*]clCreateSampler
cl_sampler clCreateSamplerWithProperties (cl_context context,
const cl_sampler_properties *sampler_properties,
cl_int *errcode_ret)
功能描述：
创建 Sampler 对象。GPU 设备通过 Sampler 读取图像时，使用更高效的纹理处理器和纹理缓存提高读取效率；使用硬件内建的边界及插值处理，降低编程复杂度。
参数说明：
- context ：为 context 对象分配 Sampler 对象。
- sampler_properties：采样器属性，包括坐标归一化，采样边界处理模式和采样插值滤波模式。
- errcode_ret ：返回错误码。errcode_ret 设为 NULL 则不再返回错误码。
- Return ：sampler 对象成功创建时，返回创建的 cl_sampler 对象，errcode_ret 返回`CL_SUCCESS`。创建失败时返回 NULL，errorcode_ret 返回错误码。Image操作实例代码
下面例子通过 OpenCV 读取图像，展示了 Image 对象的运行时 API 操作。
cl_mem memobject = {0, 0};
cv::Mat src = cv::imread(&#34;./lena512x512.jpg&#34;, 1);
cv::cvtColor(src, src, cv::COLOR_BGR2BGRA);
// create opencl memobject
cl_image_format image_format;
image_format.image_channel_order = CL_RGBA;
image_format.image_channel_data_type = CL_UNORM_INT8;
cl_image_desc image_desc;
memset(&image_desc, 0, sizeof(cl_image_desc));
image_desc.image_type = CL_MEM_OBJECT_IMAGE2D;
image_desc.image_width = src.cols;
image_desc.image_height = src.rows;
// Create Image with CL_MEM_COPY_HOST_PTR
memobject = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &image_format, &image_desc, src.data, &err_num);
memobject = clCreateImage(context, CL_MEM_WRITE_ONLY, &image_format, &image_desc, NULL, &err_num);
/**************Set Arg and Run Kenel*********************/
// copy result from device to host
cv::Mat gpu_dst(src.size(), src.type());
size_t origin = {0, 0, 0};
size_t region = {(size_t)src.cols, (size_t)src.rows, 1};
err_num = clEnqueueReadImage(command_queue, memobject, CL_TRUE, origin, region,
0, 0, gpu_dst.data, 0, NULL, NULL);
2.2.8 设置 kernel 参数
作用说明
准备好 kernel 对象和输入输出 mem 对象后，这一步设置 kernel 函数参数。
API函数说明
[*]SetKernelArg
cl_int clSetKernelArg(cl_kernel kernel,
cl_uint arg_index,
size_t arg_size,
const void *arg_value)
功能描述：
设置kernel函数的单个参数。
参数说明：
- kernel ：有效的内核对象。
- arg_index：内核函数参数索引，对于 n 个参数的内核函数，arg_index 从 0 编号到 n-1。
- arg_size ：第 arg_index 个参数占内存大小。
- arg_value：指向传入参数数据的指针。
- Return ：成功生成可执行程序返回 CL_SUCCESS，错误码参考中 5.9.2 节。实例代码
例如对于下面的 kernel 函数
// 依次设置参数
err |= clSetKernelArg(kernel, 0, sizeof(cl_mem), &image_src);
err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &image_dst);
err |= clSetKernelArg(kernel, 2, sizeof(cl_sampler), &sampler);
err |= clSetKernelArg(kernel, 3, sizeof(int), &image_width);
err |= clSetKernelArg(kernel, 4, sizeof(int), &image_height);
/* // 对应的 kernel 代码
__kernel
void kernel_func(__read_only image2d_t src_img, // arg_idx=0
__write_only image2d_t dst_img,// arg_idx=1
sampler_t sampler, // arg_idx=2
int width, // arg_idx=3
int height) // arg_idx=4
*/
2.2.9 执行内核
作用说明
设置 kernel 的任务网格尺寸，并执行内核。
API 函数说明
[*]clEnqueueNDRangeKernel
cl_int clEnqueueNDRangeKernel(cl_command_queue command_queue,
cl_kernel kernel,
cl_uint work_dim,
const size_t *global_work_offset,
const size_t *global_work_size,
const size_t *local_work_size,
cl_uint num_events_in_wait_list,
const cl_event *event_wait_list,
cl_event *event)
功能描述：
提交内核执行命令，提交后 API 立即返回 Host。设备会按照设定的 work-item 网格尺寸启动内核函数执行。
参数说明：
- command_queue ：提交内核执行命令的命令队列。
- kernel ：在设备上执行的内核函数。
- work_dim ：work-item 的组织维度，0< work_dim <= 3，全局 work-item 和工作组内work-item维度相同。
- global_work_offset ：数组，表示 0~ work_dim-1 维全局工作项 ID 的偏移量。可设为 NULL，每个维度偏移量为 0。
- global_work_size ：全局工作项尺寸数组，全局工作项总数为 global_work_size*...* global_work_size 。
- local_work_size ：工作组尺寸数组，工作组内工作项个数为local_work_size*...* local_work_size 。
- num_events_in_wait_list：event_wait_list 内事件的数目。
- event_wait_list ：cl_event 数组。执行内核操作前，需要等待 event_wait_list 内事件执行完成。如果不需要等待，event_wait_list 设为 NULL，num_events_in_wait_list 设为 0。
- event ：如果非 NULL，则此参数用于查询命令执行状态，或等待命令完成。实例代码
下面的代码按照图像数据的 width 和 height 划分全局工作项，以 16x16 个工作项为单位构建工作组。
size_t global_worksize = {width, height};
// work_group尺寸不能超过clGetKernelWorkGroupInfo查询的CL_KENREL_WORK_GROUP_SIZE
size_t local_worksize = {16, 16};
err_num = clEnqueueNDRangeKernel(command_queue, buffer_kernel, 2, NULL,
global_worksize, local_worksize, 0, NULL, NULL);
2.2.10 主机和设备同步
作用说明
OpenCL Host 和 Device 之间为异步操作，Host 读取 Device 数据前应确保 Device 端相关命令执行完成。可使用 clFinish 或 clWaitForEvents 进行同步。此外 clEnqueueRead/WriteBuffer 等也能起到阻塞同步的作用。
API函数说明
[*]clFinsh
cl_int clFinish (cl_command_queue command_queue)
功能描述：
阻塞 Host 直到 command_queue 中入队命令全部执行完成，应谨慎使用。
参数说明：
- command_queue：要阻塞等待的命令队列。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中5.15节。
[*]clWaitForEvents
cl_int clWaitForEvents (cl_uint num_events, const cl_event *event_list)
功能描述：
等待事件列表关联的设备命令执行完成。
参数说明：
- num_events：event_list 中事件数量。
- event_list ：所有要等待执行完成的事件。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中 5.11 节。
[*]clGetProfilingInfo
cl_int clGetEventProfilingInfo (cl_event event,
cl_profiling_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret)
功能描述：
使用 clEnqueue 命令返回的 Event 对象抓取设备计时。命令队列属性需要使能CL_QUEUE_PROFILING_ENABLE。
参数说明：
- event ：clEnqueue 命令返回的 Event 对象。
- param_name ：表示 Profiling 信息查询参数的枚举常量，参考表 2-9。
- param_value_size ：param_value 指向内存空间的字节数。
- param_value ：指向返回查询参数结果的内存指针。
- param_value_size_ret：返回查询参数的实际长度。
- Return ：正常执行返回 CL_SUCCESS，异常返回查询中 5.14 节。
表 2-9 Profiling 时间点查询列表
表 2-9 列出了命令执行的入队(queue)、提交(submit)、启动(start)、结束(end)和完成(complete)五个时间点，按时间线分为4个时间段。其中最值得关注的是启动到结束的时间，表示kernel函数的运行时间。
下面给出 Host 端打印各阶段时长的代码void PrintProfilingInfo(cl_event event)
{
cl_ulong t_queued;
cl_ulong t_submitted;
cl_ulong t_started;
cl_ulong t_ended;
cl_ulong t_completed;
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_QUEUED, sizeof(cl_ulong), &t_queued, NULL);
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_SUBMIT, sizeof(cl_ulong), &t_submitted, NULL);
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &t_started, NULL);
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &t_ended, NULL);
clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_COMPLETE, sizeof(cl_ulong), &t_completed, NULL);
printf(&#34;queue -> submit : %fus\n&#34;, (t_submitted - t_queued) * 1e-3);
printf(&#34;submit -> start : %fus\n&#34;, (t_started - t_submitted) * 1e-3);
printf(&#34;start -> end : %fus\n&#34;, (t_ended - t_started) * 1e-3);
printf(&#34;end -> finish : %fus\n&#34;, (t_completed - t_ended) * 1e-3);
}
对第三节示例代码的 kernel 做 Profiling 分析，可以看到 start -> end 的 kernel 函数执行时间占主体。kernel 函数内没有设备端入队指令，end->finish 耗时为 0。queue -> submit : 145.920000us
submit -> start : 88.064000us
start -> end : 10540.032000us
end -> finish : 0.000000us2.2.11 读取 Device 处理结果
作用说明
OpenCL Device kernel 以 buffer 或者 Image 的方式向 Host 返回数据，可使用 clEnqueueReadBuffer 或 clEnqueueReadImage 读取计算处理结果。
2.2.12 清理 OpenCL 资源
作用说明
OpenCL 程序执行的最后一步是手动清理在 Context 上申请的全部资源。我们建议停止使用的 OpenCL 资源尽早释放，按照与创建时相反的次序释放 OpenCL 系统资源。
OpenCL 资源普遍采用引用计数机制进行管理，当引用计数降为 0，且相关设备命令执行完成、附属资源释放后，资源对象删除。
API函数说明
[*]clReleaseMemObject
cl_int clReleaseMemObject(cl_mem memobj)
参数说明：
- memobj：减少memobj 的引用计数。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 5.5.1 节。
[*]clReleaseEvent
cl_int clReleaseEvent (cl_event event)
参数说明：
- event：减少 event 的引用计数。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 5.11 节。
[*]clReleaseSampler
cl_int clReleaseSampler (cl_sampler sampler)
参数说明：
- sampler：减少 sampler 的引用计数。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 5.7.1 节。
[*]clReleaseKernel
cl_int clReleaseKernel (cl_kernel kernel)
参数说明：
- kernel：减少 kernel 的引用计数。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 5.9.1 节。
[*]clReleaseProgram
cl_int clReleaseProgram (cl_program program)
参数说明：
- program：减少 program 的引用计数。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 5.8.1 节。
[*]clReleaseCommandQueue
cl_int clReleaseCommandQueue (cl_command_queue command_queue)
参数说明：
- command_queue：减少command_queue 的引用计数。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 5.1 节。
[*]clReleaseContext
cl_int clReleaseContext (cl_context context)
参数说明：
- context：减少 context 的引用计数。
- Return ：正常执行返回 CL_SUCCESS，异常返回值参考中 4.4 节。2.3 API 使用注意事项
OpenCL API 函数在 CPU 主机上执行，用于管理 OpenCL 资源和控制应用程序的执行。API 函数相较于设备内核代码，计算工作量更低，但不恰当地使用 API 函数可能带来很大的性能损失。开发人员可参考以下几点注意事项。
（1）内存资源的使用
[*]尽可能重用 OpenCL 内存和上下文对象，避免高开销的资源创建工作。
[*]避免在 clEnqueueNDRangeKernel 之间创建或释放 OpenCL 内存对象。
[*]由于主机和设备之间内存拷贝成本高昂，可以依据硬件厂商支持，采用不同的零拷贝机制。例如 Intel HD Graphics 使用 Map 取代 Copy，OpenCL2.0 支持的 SVM 内存，高通平台扩展的 ION 内存等等。
（2）二进制程序加载
[*]运行时加载编译 OpenCL C 源代码极为耗时。首次使用 clCreateProgramWithSource 创建 program 并编译生成后，可保存设备的可执行二进制到文件。之后使用 clCreateProgramWithBinary 加载 program。下面是同一段 OpenCL C 代码使用源码加载编译和二进制加载编译的时间对比，可以看到程序加载时间显著降低。
[*]编译好的二进制代码只适用于指定设备，如果设备加载了不匹配的二进制代码，应退回使用源码编译。
# clCreateProgramWithSource and clBuildProgram
Source Compiling consumes average time: 41441 us
# clCreateProgramWithBinary and clBuildProgram
Binary Compiling consumes average time: 39 us（3）基于事件驱动(Event-driven)设备命令运行
[*]调用阻塞式 API 会挂起主机端 CPU 进程/线程，OpenCL 程序应尽可能避免阻塞式 API 调用。
[*]OpenCL enqueue API 可以接收事件列表作为输入参数，列表中所有事件完成后，命令开始在设备端执行。使用事件列表建立命令之间的依赖关系，主机端只要将 clEnqueue 命令入队，交由设备按事件依赖执行，无需在主机端做阻塞式的同步操作。
// CL_FALSE表示 Non-blocking API，命令入队后立即返回
clEnqueueWriteBuffer(command_queue, buffer_src, CL_FALSE, 0, buffer_size_in_bytes,
host_src_matrix, 0, NULL, &write_event);
// 命令依次入队，在Device端按照event依赖执行
clEnqueueNDRangeKernel(command_queue, native_kernel, 2, NULL, global_work_size, local_work_size, 1,
&write_event, &kernel_event);
clEnqueueReadBuffer(command_queue, buffer_dst, CL_FALSE, 0, buffer_size_in_bytes,
host_dst_matrix, 1, &kernel_event, &read_event);
/*CPU do something*/
// Sync between host and device
clWaitForEvents(1, &read_event);
三、OpenCL图像转置示例
3.1 代码展示
以 8 位灰度图像转置为例编写 CPU C 代码和 OpenCL 代码，压缩包下载链接参见附录。为简化图像格式操作，用cl_uchar 矩阵表示图像。CPU 代码采用行读列写的方式，
void CpuTranspose(cl_uchar *src, cl_uchar *dst, int src_width, int src_height)
{
for (int src_row = 0; src_row < src_height; src_row++)
{
for (int src_col = 0; src_col < src_width; src_col++)
{
//执行主体
dst = src;
}
}
}
OpenCL C 源码的 kernel 函数只需要执行 CPU 循环的主体，每个工作项处理一个像素。
__kernel void TransposeKernel(__global uchar *src, __global uchar *dst, int width, int height)
{
uint g_idx = get_global_id(0);
uint g_idy = get_global_id(1);
if ((g_idx >= width) || (g_idy >= height))
return;
dst = src;
}
在 Host 端，使用本文介绍的 API，按步骤完成 OpenCL 设置，调用 clEnqueueNDRangeKernel 在设备端执行。
cl_device_id device;
cl_context context;
cl_command_queue command_queue;
cl_program program;
cl_kernel kernel;
cl_mem buffer_src;
cl_mem buffer_dst;
cl_int err_num = CL_SUCCESS;
cl_uint buffer_size_in_bytes;
timeval start;
// Step 1-3 查询平台设备并创建context
context = CreateContext(&device);
if (NULL == context)
{
printf(&#34;MainError:Create Context Failed!\n&#34;);
return -1;
}
// Step 4 创建command queue
command_queue = CreateCommandQueue(context, device);
if (NULL == command_queue)
{
printf(&#34;MainError:Create CommandQueue Failed!\n&#34;);
return -1;
}
// 读取OpenCL C源代码
char *device_source_str = ClUtilReadFileToString(&#34;kerneltest.cl&#34;);
program = CreateProgram(context, device, device_source_str);
// Step 5 创建编译program
if (NULL == program)
{
printf(&#34;MainError:Create CommandQueue Failed!\n&#34;);
return -1;
}
// Step 6 创建编译kernel
kernel = CreateKernel(program, &#34;TransposeKernel&#34;, device);
if (NULL == kernel)
{
printf(&#34;MainError:Create Kernel Failed!\n&#34;);
return -1;
}
const int c_loop_count = 30;
int width = 4096;
int height = 4096;
buffer_size_in_bytes = width * height * sizeof(cl_uchar);
cl_uchar *host_src_matrix = (cl_uchar *)malloc(buffer_size_in_bytes);
cl_uchar *host_transposed_matrix = (cl_uchar *)malloc(buffer_size_in_bytes);
cl_uchar *device_transposed_matrix = (cl_uchar *)malloc(buffer_size_in_bytes);
memset(device_transposed_matrix, 0, buffer_size_in_bytes);
DataInit(host_src_matrix, width, height);
printf(&#34;Matrix Width =%d Height=%d\n&#34;, width, height);
gettimeofday(&start, NULL);
for (int i = 0; i < c_loop_count; i++)
{
CpuTranspose(host_src_matrix, host_transposed_matrix, width, height);
}
// 计算CPU多次运行的平均时间
PrintDuration(&start, &#34;Cpu Transpose&#34;, c_loop_count);
// Step 7 创建内存对象
buffer_src = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
buffer_size_in_bytes, host_src_matrix, &err_num)
CheckClStatus(err_num, &#34;Create src buffer&#34;);
buffer_dst = clCreateBuffer(context, CL_MEM_WRITE_ONLY, buffer_size_in_bytes, NULL, &err_num);
CheckClStatus(err_num, &#34;Create dst buffer&#34;);
// Step 8 设置kernelArg
err_num = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_src);
err_num |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer_dst);
err_num |= clSetKernelArg(kernel, 2, sizeof(int), &width);
err_num |= clSetKernelArg(kernel, 3, sizeof(int), &height);
size_t global_work_size;
size_t local_work_size;
// 设置NDRange尺寸
local_work_size = 32;
local_work_size = 32;
local_work_size = 0;
global_work_size =
(width + local_work_size - 1) / local_work_size * local_work_size;
global_work_size =
(height + local_work_size - 1) / local_work_size * local_work_size;
global_work_size = 0;
printf(&#34;global_work_size=(%zu,%zu)\n&#34;, global_work_size, global_work_size);
printf(&#34;local_work_size=(%zu,%zu)\n&#34;, local_work_size, local_work_size);
cl_event kernel_event = NULL;
gettimeofday(&start, NULL);
for (int i = 0; i < c_loop_count; i++)
{
// Step 9 入队kernel执行
err_num = clEnqueueNDRangeKernel(command_queue, kernel, 2, NULL, global_work_size,
local_work_size, 0, NULL, &kernel_event);
CheckClStatus(err_num, &#34;ClEnqueueNDRangeKernel&#34;);
// Step 10 同步
err_num = clWaitForEvents(1, &kernel_event);
CheckClStatus(err_num, &#34;ClWaitForEvents&#34;);
}
// 计算GPU多次运行的平均时间
PrintDuration(&start, &#34;OpenCL Transpose&#34;, c_loop_count);
// Step 11 读取OpenCL计算结果
err_num = clEnqueueReadBuffer(command_queue, buffer_dst, CL_TRUE, 0, buffer_size_in_bytes,
device_transposed_matrix, 0, NULL, NULL);
compare(host_transposed_matrix, device_transposed_matrix, width, height);
free(device_source_str);
free(host_src_matrix);
free(host_transposed_matrix);
free(device_transposed_matrix);
// Step 12 清理OpenCL资源
clReleaseEvent(kernel_event);
clReleaseMemObject(buffer_src);
clReleaseMemObject(buffer_dst);
CleanUp(context, command_queue, program, kernel);
3.2 运行结果
在高通 865 平台，对 4096x4096 图像做转置，运行结果如下，CPU 运行时间除以 OpenCL 运行时间作为加速比，约 10.1 倍。
Cpu Transpose consume average time: 135815 us
Kernel TransposeKernel max workgroup size=1024
Kernel TransposeKernel perferred workgroup size multiple=128
global_work_size=(4096,4096)
local_work_size=(32,32)
OpenCL Transpose consume average time: 13412 us
A and B match!使用不同尺寸图像计算 OpenCL 转置加速比。随着数据量上升，OpenCL 相较于 CPU 加速优势显著。
3.3 说明
本节 OpenCL 转置示例仅为展示 OpenCL 运行时 API 的使用，程序可以深度优化并得到数倍的性能提升。优化措施包括且不限于：
四、总结
本文首先给出 OpenCL 运行时 API 的整体编程流程图，然后针对每一步介绍使用的运行时 API，讲解 API 参数，并给出编程运行实例。总结运行时 API 使用的注意事项。最后展示基于 OpenCL 的图像转置代码。在 865 平台下，对于 4096x4096 的 8 位图像加速比达到 10 倍以上。
五、工程代码
参考资料
https://www.khronos.org/registry/OpenCL/specs/opencl-2.0.pdf
Qualcomm snapdragon mobile platform opencl general programming and optimization
OpenCL 异构并行计算
Heterogeneous Computing with OpenCL 2.0
https://www.khronos.org/registry/OpenCL/sdk/2.1/docs/man/xhtml/
OpenCL in Action 期待更行opencl相关的后续内容[可怜] 查看图片 群图片能再发一次吗
页:
[1]