OpenCL demo - 1. 建立工程

很久没更新了，最近都在做OpenCL的东西，打算分享一下这方面的基础知识。因为现在并行计算比较普遍，区块链和深度学习都需要做大量的计算，所以OpenCL和Cuda的优势开始体现出来了。但是确实不知道OpenCL能撑多久。。。不过概念搞熟之后，cuda和OpenCL都是差不多的。

以后的分享还是以代码为主，写代码的过程讲一些我理解的概念~~

因为OpenCL在桌面机用的多，所以代码主要是linux的。windows暂时不会去弄，环境应该都比较好配置。移动端的话...额....，手上没设备。。。先算了吧。。。

下面我们直接开始写代码吧。OpenCL的头文件很简单，直接

#include <CL/opencl.h>

然后我们先获取platform id。 platform可以看做是一个设备集合的载体吧。

cl_platform_id platform_id;
checkErr(clGetPlatformIDs(1, &platform_id, nullptr));

然后通过这个platform_id, 我们就可以获取设备了。设备可以看做是所有计算单元的集合。

cl_device_id device_id;
checkErr(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, 0));

通过设备创建一个context。context中含有我们需要运行的指令的所有资源和状态。

cl_context context = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, nullptr);

然后创建command queue。command queue是真正显卡上运行的指令的command，它属于context。我们让显卡执行的所有东西都会保存在command queue上面。我们也可以同时创建多个command queue来同时执行。这个我们以后如果讨论async的话可以单独讲一下。

cl_command_queue queue = clCreateCommandQueue(context, device_id, 0, nullptr);

到这里位置，我们需要的设备相关的东西就创建完成了。现在开始创建我们显卡执行的代码。首先是一个kernel的源代码和创建并编译一个program 程序。program程序中可以包含一个或者多个kernel，也可以包含其他的函数。kernel是一个真正的执行单元，也可以认为是一个显卡的入口函数。

const char * src =
  "__kernel void demo1() {\n"
  "  int a = 0;\n"
  "}";

  cl_program program = clCreateProgramWithSource(context, 1, &src, nullptr, nullptr);
  checkErr(clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr));

从上面的src可以看到，我们这个kernel其实什么事情也没做。kernel函数必须包含关键字__kernel来向openCL表明这是一个kernel函数。其他的语法基本和C语言类似。创建并编译好一个program之后，我们就可以从这个program中创建kernel了，创建kernel的时候需要指明kernel的名字，在这里是demo1.

cl_kernel kernel = clCreateKernel(program, "demo1", nullptr);

创建好kernel之后，我们就可以让显卡来运行这个kernel了。这里指明了global size和local size，下次具体讨论这个size的含义，这里将两个都设为1，其实就是只运行一下kernel。这里我们还指明了command queue，表明我们需要把这个执行放在哪个command queue上面。同时finish也需要指定是finish哪个command queue。

size_t global_size[] = {1};
size_t local_size[] = {1};
checkErr(clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, global_size, local_size, 0, nullptr, nullptr));
clFinish(queue);

 

到此为止，一个非常非常简单的opencl程序就完成了。虽然什么事情也没有做，但是这个流程已经完整了，并且可以编译和运行。后面我们会逐步细化并讨论更多opencl的特性。

代码用了cmake，用法可以查下网上。代码上传到：https://download.csdn.net/download/hoytgm/10553847