看了mnn、mace和ncnn的源码,对于深度学习手机gpu计算的优化,mace选择了opencl,而ncnn选择了vulkan,mnn选择了opencl、vulkan、metal、opengl。那本篇就来介绍一下嵌入式(手机)下opencl的异构编程。下篇vulkan。
每年看各手机厂商发布会的时候,都会提到该手机除了牛逼的某某处理器之外,还配备了某某协处理器,其实这个协处理器就相当于手机上的GPU(图形处理单元)。除了手机,还有其他的一些嵌入式设备也配备了其他类型的协处理器,比如Xilinx的Zynq用FPGA来补充ARM Cortex-A9处理器的不足,再比如Texas Instruments的Keystone II用DSP补充了ARM Cortex A15的不足。这些协处理器都支持通用编程使用开发计算语言(OpenCL)。说回到手机,现在主流的手机上用的GPU主要是以下两种:1)ARM Mali GPU;2)Qualcomm Adreno GPU。这俩都支持opencl。
- OpenCL规范
OpenCL规范由四个模型组成,总结如下:
- 平台模型:描述了协同执行的单个处理器(宿主机)及一个或多个能执行OpenCL代码的处理器(设备)。它定义了一个抽象的硬件模型,供编程人员用于编写能够再这些设备上执行的OpenCL C函数(称作kernel)。
- 执行模型:定义了在主机上如何配置OpenCL环境以及如何在设备上执行kernel。这包括在主机端建立OpenCL上下文,提供主机-设备之间的交互机制,定义一个并发模型供在设备上执行kernel所用。
- 内存模型: 定义被kernel所用的抽象内存层次,无需考虑实际的底层内存架构。尽管内存模型十分接近当前的GPU内存层次,但同样也适用于其他硬件加速器。
- 编程模型:定义了如何将并发模型映射到物理硬件上。
- OpenCL应用步骤
一个基于异构平台的应用一般会包含以下步骤:
- 找出组成异构平台的所有组件
- 考察组件的特征,这样就能使得软件根据不同的硬件特征来实现。
- 创建在平台上运行的一组内核。
- 设置与计算相关的存储对象
- 在合适的组件上以正确的顺序执行内核。
- 收集结果。
-
OpenCL编程
OpenCl的编程框架组成包括:
平台API:平台API定义了宿主机程序发现OpenCL设备所用的函数以及这些函数的功能,另外还定义为OpenCL应用创建上下文(上下文表示的是程序运行时所拥有的所有软硬件资源+内存+处理器)的函数。这里的平台指的是宿主机、OpenCL设备和OpenCL框架的组合。
运行时API:平台API主要用来创建上下文,运行时API则强调使用这个上下文满足应用需求的函数集,用来管理上下文来创建命令队列以及运行时发生的其它操作。例如,将命令提交到命令队列的函数。
编程语言:用来编写内核代码的编程语言,基于ISO C99标准的一个扩展子集,通常称为OpenCL C编程语言。
- OpenCL的编程步骤如下:
- Discover and initialize the platforms
调用两次clGetPlatformIDs函数,第一次获取可用的平台数量,第二次获取一个可用的平台。 - Discover and initialize the devices
调用两次clGetDeviceIDs函数,第一次获取可用的设备数量,第二次获取一个可用的设备。 - Create a context(调用clCreateContext函数)
上下文context可能会管理多个设备device。 - Create a command queue(调用clCreateCommandQueue函数)
一个设备device对应一个command queue。
上下文conetxt将命令发送到设备对应的command queue,设备就可以执行命令队列里的命令。 - Create device buffers(调用clCreateBuffer函数)
Buffer中保存的是数据对象,就是设备执行程序需要的数据保存在其中。
Buffer由上下文conetxt创建,这样上下文管理的多个设备就会共享Buffer中的数据。 - Write host data to device buffers(调用clEnqueueWriteBuffer函数)
- Create and compile the program
创建程序对象,程序对象就代表你的程序源文件或者二进制代码数据。 - Create the kernel(调用clCreateKernel函数)
根据你的程序对象,生成kernel对象,表示设备程序的入口。 - Set the kernel arguments(调用clSetKernelArg函数)
- Configure the work-item structure(设置worksize)
配置work-item的组织形式(维数,group组成等) - Enqueue the kernel for execution(调用clEnqueueNDRangeKernel函数)
将kernel对象,以及 work-item参数放入命令队列中进行执行。 - Read the output buffer back to the host(调用clEnqueueReadBuffer函数)
- Release OpenCL resources(至此结束整个运行过程)
- OpenCL实战之矩阵相乘
了解了OpenCL的的理论知识,这里演示一个MP的矩阵和PN的矩阵相乘的例子。
首先我们知道我们的程序分为两部分,一部分是我们的kernel运行再gpu,一部分是我们的主程序运行再cpu。kernel可以这样写:
__kernel void matrix_mult(
const int Ndim,
const int Mdim,
const int Pdim,
__global const float* A,
__global const float* B,
__global float* C)
{
int i = get_global_id(0);
int j = get_global_id(1);
int k;
float tmp;
if ((i < Ndim) && (j < Mdim)) {
tmp = 0.0;
for (k = 0; k < Pdim; k++)
tmp += A[i*Pdim + k] * B[k*Mdim + j];
C[i*Mdim + j] = tmp;
}
}
那么主程序就可以按照上面的理论步骤来写:
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <iostream>
#include <fstream>
using namespace std;
#define NWITEMS 6
#pragma comment (lib,"OpenCL.lib")
//把文本文件读入一个 string 中
int convertToString(const char *filename, std::string& s)
{
size_t size;
char* str;
std::fstream f(filename, (std::fstream::in | std::fstream::binary));
if (f.is_open())
{
size_t fileSize;
f.seekg(0, std::fstream::end);
size = fileSize = (size_t)f.tellg();
f.seekg(0, std::fstream::beg);
str = new char[size + 1];
if (!str)
{
f.close();
return NULL;
}
f.read(str, fileSize);
f.close();
str[size] = '\0';
s = str;
delete[] str;
return 0;
}
printf("Error: Failed to open file %s\n", filename);
return 1;
}
int main()
{
cl_uint status;
cl_platform_id platform;
//创建平台对象
status = clGetPlatformIDs(1, &platform, NULL);
cl_device_id device;
//创建 GPU 设备
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU,
1,
&device,
NULL);
//创建context
cl_context context = clCreateContext(NULL,
1,
&device,
NULL, NULL, NULL);
//创建命令队列
cl_command_queue commandQueue = clCreateCommandQueue(context,
device,
CL_QUEUE_PROFILING_ENABLE, NULL);
if (commandQueue == NULL)
perror("Failed to create commandQueue for device 0.");
//建立要传入从机的数据
/******** 创建内核和内存对象 ********/
const int Ndim = 4;
const int Mdim = 5;
const int Pdim = 3;
int szA = Ndim * Pdim;
int szB = Pdim * Mdim;
int szC = Pdim * Mdim;
float *A;
float *B;
float *C;
A = (float *)malloc(szA * sizeof(float));
B = (float *)malloc(szB * sizeof(float));
C = (float *)malloc(szC * sizeof(float));
int i, j;
for (i = 0; i < szA; i++)
A[i] = (float)((float)i + 1.0);
for (i = 0; i < szB; i++)
B[i] = (float)((float)i + 1.0);
//创建三个 OpenCL 内存对象,并把buf1 的内容通过隐式拷贝的方式
//拷贝到clbuf1, buf2 的内容通过显示拷贝的方式拷贝到clbuf2
cl_mem memObjects[3] = { 0, 0, 0 };
memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(float)* szA, A, NULL);
memObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(float)* szB, B, NULL);
memObjects[2] = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(float)* szC, C, NULL);
if (memObjects[0] == NULL || memObjects[1] == NULL ||memObjects[2] == NULL)
perror("Error in clCreateBuffer.\n");
const char * filename = "Vadd.cl";
std::string sourceStr;
status = convertToString(filename, sourceStr);
if (status)
cout << status << " !!!!!!!!" << endl;
const char * source = sourceStr.c_str();
size_t sourceSize[] = { strlen(source) };
//创建程序对象
cl_program program = clCreateProgramWithSource(
context,
1,
&source,
sourceSize,
NULL);
//编译程序对象
status = clBuildProgram(program, 1, &device, NULL, NULL, NULL);
if (status)
cout << status << " !!!!!!!!" <<endl;
if (status != 0)
{
printf("clBuild failed:%d\n", status);
char tbuf[0x10000];
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf,
NULL);
printf("\n%s\n", tbuf);
//return −1;
}
//创建 Kernel 对象
cl_kernel kernel = clCreateKernel(program, "matrix_mult", NULL);
//设置 Kernel 参数
cl_int clnum = NWITEMS;
status = clSetKernelArg(kernel, 0, sizeof(int), &Ndim);
status = clSetKernelArg(kernel, 1, sizeof(int), &Mdim);
status = clSetKernelArg(kernel, 2, sizeof(int), &Pdim);
status = clSetKernelArg(kernel, 3, sizeof(cl_mem), &memObjects[0]);
status = clSetKernelArg(kernel, 4, sizeof(cl_mem), &memObjects[1]);
status = clSetKernelArg(kernel, 5, sizeof(cl_mem), &memObjects[2]);
if (status)
cout << "参数设置错误" << endl;
//执行 kernel
size_t global[2];
cl_event prof_event;
cl_ulong ev_start_time = (cl_ulong)0;
cl_ulong ev_end_time = (cl_ulong)0;
double rum_time;
global[0] = (size_t)Ndim;
global[1] = (size_t)Mdim;
status = clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL,
global, NULL, 0, NULL, &prof_event);
if (status)
cout << "执行内核时错误" << endl;
clFinish(commandQueue);
//读取时间
status = clGetEventProfilingInfo(prof_event,CL_PROFILING_COMMAND_QUEUED,
sizeof(cl_ulong),&ev_start_time,NULL);
status = clGetEventProfilingInfo(prof_event,CL_PROFILING_COMMAND_END,
sizeof(cl_ulong),&ev_end_time,NULL);
if (status)
perror("读取时间的时候发生错误\n");
rum_time = (double)(ev_end_time - ev_start_time);
cout << "执行时间为:" << rum_time << endl;
//数据拷回 host 内存
status = clEnqueueReadBuffer(commandQueue, memObjects[2],CL_TRUE, 0,
sizeof(float)* szC, C,0, NULL, NULL);
if (status)
perror("读回数据的时候发生错误\n");
//结果显示
printf("\nArray A:\n");
for (i = 0; i < Ndim; i++) {
for (j = 0; j < Pdim; j++)
printf("%.3f\t", A[i*Pdim + j]);
printf("\n");
}
printf("\nArray B:\n");
for (i = 0; i < Pdim; i++) {
for (j = 0; j < Mdim; j++)
printf("%.3f\t", B[i*Mdim + j]);
printf("\n");
}
printf("\nArray C:\n");
for (i = 0; i < Ndim; i++) {
for (j = 0; j < Mdim; j++)
printf("%.3f\t", C[i*Mdim + j]);
printf("\n");
}
cout << endl;
if (A)
free(A);
if (B)
free(B);
if (C)
free(C);
//删除 OpenCL 资源对象
clReleaseMemObject(memObjects[2]);
clReleaseMemObject(memObjects[1]);
clReleaseMemObject(memObjects[0]);
clReleaseProgram(program);
clReleaseCommandQueue(commandQueue);
clReleaseContext(context);
system("pause");
return 0;
}
参考链接:
https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html
http://www.cnblogs.com/xudong-bupt/p/3582780.html
https://blog.csdn.net/c602273091/article/details/45418129