Read the fucking official documents!
--By 鲁迅A picture is worth a thousand words.
--By 高尔基说明:
OpenCL
,目标是浅显易懂,如果没有达到这个效果,就当我没说这话;Middleware
的系统软件工程师,不是一个好码农;kernels
的语言(基于C99);opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz
以人工智能场景为例来理解一下,假如在某个AI芯片上跑人脸识别应用,CPU擅长控制,AI processor擅长计算,软件的flow就可以进行拆分,用CPU来负责控制视频流输入输出前后处理,AI processor来完成深度学习模型运算完成识别,这就是一个典型的异构处理场景,如果该AI芯片的SDK支持OpenCL,那么上层的软件就可以基于OpenCL进行开发了。
话不多说,看看OpenCL的架构吧。
OpenCL架构,可以从平台模型、内存模型、执行模型、编程模型四个角度来展开。
平台模型:硬件拓扑关系的抽象描述
Compute Unit(CU)
;Processing Unit(PE)
,最终的计算由PE来完成;执行模型:Host如何利用OpenCL Device的计算资源完成高效的计算处理过程
OpenCL的Execution Model由两个不同的执行单元定义:1)运行在OpenCL设备上的kernel;2)运行在Host上的Host program;其中,OpenCL使用Context代表kernel的执行环境:
Context包含以下资源:
有两种方式来找到work-item:
以一维为例:
以二维为例:
三维的方式也类似,略去。
内存模型:Host和OpenCL Device怎么来看待数据
OpenCL的内存模型中,包含以下几类类型的内存:
下边来一个实际的代码测试跑跑,Talk is cheap, show me the code!
opencl_runtime_16.1.2_x64_rh_6.4.0.37.tgz
);#include <iostream>
#include <fstream>
#include <sstream>
#include <CL/cl.h>
const int DATA_SIZE = 10;
int main(void)
{
/* 1. get platform & device information */
cl_uint num_platforms;
cl_platform_id first_platform_id;
clGetPlatformIDs(1, &first_platform_id, &num_platforms);
/* 2. create context */
cl_int err_num;
cl_context context = nullptr;
cl_context_properties context_prop[] = {
CL_CONTEXT_PLATFORM,
(cl_context_properties)first_platform_id,
0
};
context = clCreateContextFromType(context_prop, CL_DEVICE_TYPE_CPU, nullptr, nullptr, &err_num);
/* 3. create command queue */
cl_command_queue command_queue;
cl_device_id *devices;
size_t device_buffer_size = -1;
clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, nullptr, &device_buffer_size);
devices = new cl_device_id[device_buffer_size / sizeof(cl_device_id)];
clGetContextInfo(context, CL_CONTEXT_DEVICES, device_buffer_size, devices, nullptr);
command_queue = clCreateCommandQueueWithProperties(context, devices[0], nullptr, nullptr);
delete [] devices;
/* 4. create program */
std::ifstream kernel_file("vector_add.cl", std::ios::in);
std::ostringstream oss;
oss << kernel_file.rdbuf();
std::string srcStdStr = oss.str();
const char *srcStr = srcStdStr.c_str();
cl_program program;
program = clCreateProgramWithSource(context, 1, (const char **)&srcStr, nullptr, nullptr);
/* 5. build program */
clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr);
/* 6. create kernel */
cl_kernel kernel;
kernel = clCreateKernel(program, "vector_add", nullptr);
/* 7. set input data && create memory object */
float output[DATA_SIZE];
float input_x[DATA_SIZE];
float input_y[DATA_SIZE];
for (int i = 0; i < DATA_SIZE; i++) {
input_x[i] = (float)i;
input_y[i] = (float)(2 * i);
}
cl_mem mem_object_x;
cl_mem mem_object_y;
cl_mem mem_object_output;
mem_object_x = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_x, nullptr);
mem_object_y = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float) * DATA_SIZE, input_y, nullptr);
mem_object_output = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * DATA_SIZE, nullptr, nullptr);
/* 8. set kernel argument */
clSetKernelArg(kernel, 0, sizeof(cl_mem), &mem_object_x);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &mem_object_y);
clSetKernelArg(kernel, 2, sizeof(cl_mem), &mem_object_output);
/* 9. send kernel to execute */
size_t globalWorkSize[1] = {DATA_SIZE};
size_t localWorkSize[1] = {1};
clEnqueueNDRangeKernel(command_queue, kernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
/* 10. read data from output */
clEnqueueReadBuffer(command_queue, mem_object_output, CL_TRUE, 0, DATA_SIZE * sizeof(float), output, 0, nullptr, nullptr);
for (int i = 0; i < DATA_SIZE; i++) {
std::cout << output[i] << " ";
}
std::cout << std::endl;
/* 11. clean up */
clRetainMemObject(mem_object_x);
clRetainMemObject(mem_object_y);
clRetainMemObject(mem_object_output);
clReleaseCommandQueue(command_queue);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseContext(context);
return 0;
}
vector_add.cl
文件中内容如下:
__kernel void vector_add(__global const float *input_x,
__global const float *input_y,
__global float *output)
{
int gid = get_global_id(0);
output[gid] = input_x[gid] + input_y[gid];
}
The OpenCL Specification
如果对你有用的话,在看,分享,打赏三连吧,谢谢。