180d59932Sopenharmony_ci# OpenCL<sup>TM</sup> API Headers
280d59932Sopenharmony_ci
380d59932Sopenharmony_ci仓库包含C语言的OpenCL API。OpenCL扩展了GPU用于图形渲染之外的能力,将通用计算并行化。OpenCL可用于图像处理、AI、高性能计算等场景的加速。推荐在计算密集型的任务以及可以并行计算的场景使用OpenCL。
480d59932Sopenharmony_ci
580d59932Sopenharmony_ci例如在图像处理场景,CPU进行编解码、滤镜等特效处理较慢,使用OpenCL有十倍的加速效果。
680d59932Sopenharmony_ci
780d59932Sopenharmony_ci例如在AI推理场景,使用OpenCL可以将推理性能提高一倍。
880d59932Sopenharmony_ci
980d59932Sopenharmony_ciOpenHarmony引入后,新增了一层封装层,动态链接查找OpenCL库,以此进行解耦。目前仅支持native层调用OpenCL。
1080d59932Sopenharmony_ci
1180d59932Sopenharmony_ci## 目录结构
1280d59932Sopenharmony_ci
1380d59932Sopenharmony_ci```
1480d59932Sopenharmony_ciREADME.md               英文说明
1580d59932Sopenharmony_ciREADME_zh.md            中文说明
1680d59932Sopenharmony_ciLICENSE                 证书文件
1780d59932Sopenharmony_ciCL/                     原CL头文件
1880d59932Sopenharmony_ciinclude/                封装层CL头文件
1980d59932Sopenharmony_cisrc/                    封装层CL实现
2080d59932Sopenharmony_ci```
2180d59932Sopenharmony_ci
2280d59932Sopenharmony_ci## OpenHarmony如何集成OpenCL
2380d59932Sopenharmony_ci### 1.头文件引入
2480d59932Sopenharmony_ci```c
2580d59932Sopenharmony_ci#define USE_OPENCL_WRAPPER
2680d59932Sopenharmony_ci#include "opencl_wrapper.h"
2780d59932Sopenharmony_ci```
2880d59932Sopenharmony_ci### 2.BUILD.gn添加引用
2980d59932Sopenharmony_ci```c
3080d59932Sopenharmony_cideps += ["//third_party/openCL:libcl"]
3180d59932Sopenharmony_ci```
3280d59932Sopenharmony_ci### 3.调用OpenCL函数过程举例
3380d59932Sopenharmony_ci```c
3480d59932Sopenharmony_ci// 准备OpenCL的kernel程序
3580d59932Sopenharmony_ciconst char* program_source =
3680d59932Sopenharmony_ci      "__kernel void test_main(read_only image2d_t inputImage) {\n"
3780d59932Sopenharmony_ci      "  ...\n"
3880d59932Sopenharmony_ci      "}";
3980d59932Sopenharmony_ci```
4080d59932Sopenharmony_ci
4180d59932Sopenharmony_ci```c
4280d59932Sopenharmony_ci// 通过wrapper层的函数判断是否cl初始化成功
4380d59932Sopenharmony_ci// 如果返回false,说明没有实际的驱动
4480d59932Sopenharmony_cibool cl_ok = OHOS::InitOpenCL();
4580d59932Sopenharmony_ci
4680d59932Sopenharmony_ci// 获取设备信息
4780d59932Sopenharmony_cicl_int err;
4880d59932Sopenharmony_cicl_platform_id platform_id;
4980d59932Sopenharmony_cicl_device_id device_id;
5080d59932Sopenharmony_ciclGetPlatformIDs(1, &platform_id, NULL);
5180d59932Sopenharmony_ciclGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
5280d59932Sopenharmony_ci// 创建OpenCL上下文
5380d59932Sopenharmony_cicontext_ = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
5480d59932Sopenharmony_ciqueue_ = clCreateCommandQueueWithProperties(context_, device_id, 0, &err);
5580d59932Sopenharmony_ci// 使用上述源码创建OpenCL程序
5680d59932Sopenharmony_cicl_program p = clCreateProgramWithSource(context, 1, &program_source, nullptr, nullptr);
5780d59932Sopenharmony_ci// 创建kernel程序,对应上述源码中的函数名
5880d59932Sopenharmony_cikernel_ = clCreateKernel(program, "test_main", &err);
5980d59932Sopenharmony_ci// 创建OpenCL可以识别的图片
6080d59932Sopenharmony_cicl_image_format image_format;
6180d59932Sopenharmony_ciimage_format.image_channel_order = CL_RGBA;
6280d59932Sopenharmony_ciimage_format.image_channel_data_type = CL_UNORM_INT8;
6380d59932Sopenharmony_cicl_image_desc desc = {CL_MEM_OBJECT_IMAGE2D, width, height};
6480d59932Sopenharmony_cicl_mem input_image = clCreateImage(context_, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, image_format, &desc, const_cast<void*(pixmap.addr()), &err);
6580d59932Sopenharmony_ci// 将该图片传入kernel程序
6680d59932Sopenharmony_cierr |= clSetKernelArg(kernel_, 0, sizeof(cl_mem), &input_image);
6780d59932Sopenharmony_ci// 在GPU上执行OpenCL程序
6880d59932Sopenharmony_cierr = clEnqueueNDRangeKernel(queue_, kernel_, 2, NULL, global,local, 0, NULL, NULL);
6980d59932Sopenharmony_ci// 等待执行结束
7080d59932Sopenharmony_ciclFinish(queue_);
7180d59932Sopenharmony_ci// 读取执行结果,从GPU传输回CPU
7280d59932Sopenharmony_cierr = clEnqueueReadBuffer(queue_, cl_errs, CL_TRUE, 0, pixmap.size(), errs, 0, NULL, NULL);
7380d59932Sopenharmony_ci// 使用完毕需要释放cl内存
7480d59932Sopenharmony_ciclReleaseMemObject(input_image);
7580d59932Sopenharmony_ci// 程序不再需要执行的时候需要释放GPU资源(OpenCL上下文)
7680d59932Sopenharmony_ciclReleaseKernel(kernel_);
7780d59932Sopenharmony_ciclReleaseCommandQueue(queue_);
7880d59932Sopenharmony_ciclReleaseContext(context_);
7980d59932Sopenharmony_ci```
8080d59932Sopenharmony_ci
8180d59932Sopenharmony_ci## OpenCL使用文档
8280d59932Sopenharmony_ci
8380d59932Sopenharmony_ciAPI官方文档  https://registry.khronos.org/OpenCL/
8480d59932Sopenharmony_ci
8580d59932Sopenharmony_ciAPI详细定义 https://registry.khronos.org/OpenCL/sdk/3.0/docs/man/html/
8680d59932Sopenharmony_ci
8780d59932Sopenharmony_ciAPI一张图的使用说明 https://www.khronos.org/files/opencl30-reference-guide.pdf
8880d59932Sopenharmony_ci
8980d59932Sopenharmony_ci## License
9080d59932Sopenharmony_ci
9180d59932Sopenharmony_ci见 [LICENSE](LICENSE).
9280d59932Sopenharmony_ci
9380d59932Sopenharmony_ci---
9480d59932Sopenharmony_ci
9580d59932Sopenharmony_ciOpenCL and the OpenCL logo are trademarks of Apple Inc. used by permission by Khronos.
96