opencl安装_OpenCL编程简介: 第一个程序

随着深度学习的流行, 并行计算开始进入更多人的视野, 所以打算在知乎连续更新关于并行计算的编程, 其中主要是OpenCL和Cuda. OpenCL是开放标准, 有Spec可以参照, 所以先讲OpenCL, 然后再引出更易用的Cuda. 文章主要是以代码为基础, 然后讲解各个方面的概念, 这样我觉得更容易上手.

OpenCL和Cuda的安装在这里就不用介绍了, 毕竟NV的显卡装上驱动就有这个了. 其他家的显卡装上则只有OpenCL.

我们要实现一个完整的OpenCL小程序 (除了一些释放的操作, ), 这个小程序帮助我们计算A+B=C. A和B都是1 ~ 10000递增的整数. 我们这个小程序可以直接编译和执行. 下面我们就开始讲解每部分代码和一些基本概念. (上面是代码, 下面是讲解)

#include <memory>
#include <iostream>
#include <cstring>
#include <CL/opencl.h>

int main(int argc, char** argv) {
  cl_int err;
  cl_platform_id platform;
  err = clGetPlatformIDs(1, &platform, nullptr);
  if (err != CL_SUCCESS) {
    std::cout << "Cannot get platform" << std::endl;
    return -1;
  }

  cl_device_id device;
  err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 1, &device, nullptr);
  if (err != CL_SUCCESS) {
    std::cout << "Cannot get device" << std::endl;
    return -1;
  }

最开始我们需要获取一个OpenCL的platform. platform是提供一个设备合集的平台, 我们所有运行的命令和资源都是在这个平台里面进行的. 获取到platform之后, 就可以在这个platform下面获取一个设备device. device是一系列计算单元的合集, 并且我们所有的命令都是最终下发到这个设备上面执行的.

  cl_context context = clCreateContext(nullptr, 1, &device, nullptr, 
                                       nullptr, &err);
  if (err != CL_SUCCESS) {
    std::cout << "Create context failed" << std::endl;
    return -1;
  }

  cl_command_queue queue = clCreateCommandQueue(context, device, 0, &err);
  if (err != CL_SUCCESS) {
    std::cout << "Create command queue failed" << std::endl;
    return -1;
  }

从device上面就可以开始创建上下文context了. context是一个包含了所有kernel, 资源以及命令的环境. 所以我们需要创建kernel, 资源以及command的时候, 都需要context的参与. 有了context之后, 就可以开始创建command queue了. command queue是一个用来存储特定设备上面的command的队列, 这些command包含计算, 读写内存等.

  const uint32_t cal_num = 10000;
  uint32_t* hA = new uint32_t[cal_num];
  uint32_t* hB = new uint32_t[cal_num];
  uint32_t* hC = new uint32_t[cal_num];

  // initialize data
  memset(hC, 0, sizeof(uint32_t) * cal_num);
  for (uint32_t i = 0; i < cal_num; i++) {
    hA[i] = hB[i] = i;
  }

  cl_mem mA = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                             sizeof(uint32_t) * cal_num, hA, nullptr);
  cl_mem mB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                             sizeof(uint32_t) * cal_num, hB, nullptr);
  cl_mem mC = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(uint32_t) * cal_num,
                             nullptr, nullptr);
  if (mA == nullptr || mB == nullptr || mC == nullptr) {
    std::cout << "Create buffer failed" << std::endl;
    return -1;
  }

接下来是用C++创建三块内存hA, hB和hC. hA和hB初始化成0 ~ (10000 - 1), hC初始化成0. 然后就是创建三个OpenCL的内存mA, mB和mC. 其中mA和mB在创建的时候指定了一个flag叫做CL_MEM_COPY_HOST_PTR, 并且把hA和hB作为参数传了进去. 这个意思就是说, 我们在创建OpenCL的内存mA和mB的时候, 需要把mA和mB的值初始化成hA和hB的值.

  const char* program_source =
      "__kernel void test_main(__global const uint* A, __global const uint* B, __global uint* C) {n"
      "  size_t idx = get_global_id(0);n"
      "  C[idx] = A[idx] + B[idx];n"
      "}";

数据准备好了, 就开始创建执行的代码. 这里我们写了一个内核函数program_source. 内核函数也叫做kernel. program_source里面我们先只关心函数体里面的那两行代码. 第一行是获取一个全局的id. 因为我们要计算10000次, 所以这个全局id就是告诉我们当前执行的是10000次的哪一次. 第二行代码比较简单, 就是拿这个id去寻址A, B和C, 然后把A+B的值赋给C. 这里的A, B和C其实就是我们刚才创建的mA, mB和mC, 但是他们之间还需要一个对应关系来连接, 这个我们马上讲.

  cl_program program = clCreateProgramWithSource(context, 1, &program_source,
                                                 nullptr, nullptr);
  if (program == nullptr) {
    std::cout << "Create program failed" << std::endl;
    return -1;
  }

  err = clBuildProgram(program, 0, nullptr, nullptr, nullptr, nullptr);
  if (err != CL_SUCCESS) {
    std::cout << "Build program failed" << std::endl;
    return -1;
  }

  cl_kernel kernel = clCreateKernel(program, "test_main", nullptr);
  if (kernel == nullptr) {
    std::cout << "Create kernel failed" << std::endl;
    return -1;
  }

program_source有了之后, 就是创建一个program. program是一系列的kernel的合集, 这里我们只有一个kernel. 同时创建program用的是字符串代码, 也就是source, 所以我们调用的是clCreateProgramWithSource. 有了program之后, 需要OpenCL提供的编译器来编译这个program的内容, 也就是build program. build成功之后, 我们就可以从这个program里面创建我们需要执行的内核函数kernel了. 创建kernel的时候需要提供一个函数名, 也就是test_main.

  err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &mA);
  err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &mB);
  err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &mC);
  if (err != CL_SUCCESS) {
    std::cout << "Set kernel arg failed" << std::endl;
    return -1;
  }

kernel创建完成之后, 就需要指定我们刚才说的mA, mB, mC和A, B, C的对应关系了. 这个对应关系通过clSetKernelArg来完成. SetKernelArg的几个参数表明了要把哪个buffer对应到哪个kernel的第几个参数上面. 我么这个程序里面一共三个参数, 序号就是0, 1, 2. 所以我们需要调用三次, 分别把mA指定到kernel的0, mB指定到kernel的1, mC指定到kernel的2.

  size_t global_size[] {cal_num};
  size_t local_size[] {cal_num / 10};
  err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, global_size,
                               local_size, 0, nullptr, nullptr);
  if (err != CL_SUCCESS) {
    std::cout << "Run kernel failed" << std::endl;
    return -1;
  }

到此为止, 我们的准备工作就已经做完了, 可以开始执行了. 在执行之前, 我们首先需要指定我们总共要执行多少次, 也就是global_size, 在这里是10000次. 同时, 这个值也是决定了kernel里面我们获取到的那个全局的id最大是10000-1. 其实要指定一个local_size, local_size可以先理解成我们每执行一组要执行多少次, 这里是1000次. 从local_size和global_size可以体现出, 我们需要在计算单元上面执行10组, 每组执行1000次, 总共就是10000次. 最后就是调用clEnqueueNDRangeKernel, 告诉OpenCL我们需要执行的kernel和参数.

  err = clEnqueueReadBuffer(queue, mC, CL_TRUE, 0, sizeof(uint32_t) * cal_num,
                            hC, 0, nullptr, nullptr);
  if (err != CL_SUCCESS) {
    std::cout << "Read data failed" << std::endl;
    return -1;
  }

  // check one output data
  if (hC[1024] != hA[1024] + hB[1024]) {
    std::cout << "Data calculation failed" << std::endl;
    return -1;
  }

  return 0;
}

执行完计算之后, 我们怎么把结果拿到呢? 我们知道mC里面就有我们要的结果, 但是mC是不能直接索引来读取的, 还需要一个函数叫做clEnqueueReadBuffer来完成, 这个函数就告诉了OpenCL我们要把mC的数据读取之后放在哪里, 也就是放在hC里面. 完成之后, hC里面就有了最终的结果, 可以用来继续做后续的工作了.

这样一个简单的OpenCL程序就完成了. 后面我们要对这个程序进行展开, 详细的讲解一下每个部分以及更多的概念. 同时, 最后附上一个CMakeList.txt, 方便大家编译和执行.

add_definitions("-DSRCDIR="${CMAKE_CURRENT_SOURCE_DIR}"")
include_directories(/usr/local/cuda-10.0/include)

add_executable("test" main.cc)

target_compile_options("test" PRIVATE -std=c++11)
target_link_libraries("test" OpenCL)