OpenCL并行编程

看了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编程

20181105142838114.png

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