OpenCL程序编程基本流程

来源：互联网发布：港湾网络高层编辑：程序博客网时间：2024/05/29 19:50

本文转自http://www.photoneray.com/opencl_01/#program

OpenCL作为一门开源的异构并行计算语言，设计之初就是使用一种模型来模糊各种硬件差异。作为软件开发人员，我们关注的就是它的编程模型。OpenCL程序的流程大致如下：

Platform
- 查询并选择一个 platform
- 在 platform 上创建 context
- 在 context 上查询并选择一个或多个 device
Running time
- 加载 OpenCL 内核程序并创建一个 program 对象
- 为指定的 device 编译 program 中的 kernel
- 创建指定名字的 kernel 对象
- 为 kernel 创建内存对象
- 为 kernel 设置参数
- 在指定的 device 上创建 command queue
- 将要执行的 kernel 放入 command queue
- 将结果读回 host
资源回收
总结
示例代码

下面我们通过一个具体的示例程序来说明这些步骤。

使用 OpenCL API 编程与一般 C/C++ 引入第三方库编程没什么区别。所以，首先要做的自然是 include 相关的头文件。由于在 MacOS X 10.6下OpenCL的头文件命名与其他系统不同，通常使用一个#if defined进行区分，代码如下：

#if defined(__APPLE__) || defined(__MACOSX)#include <OpenCL/cl.hpp>#else#include <CL/cl.h>#endif

接下来我们就进入真正的编码流程了。

Platform

### 查询并选择一个 platform ###首先我们要取得系统中所有的 OpenCL platform。所谓的 platform 指的就是硬件厂商提供的 OpenCL 框架，不同的 CPU/GPU 开发商（比如 Intel、AMD、Nvdia）可以在一个系统上分别定义自己的 OpenCL 框架。所以我们需要查询系统中可用的 OpenCL 框架，即 platform。使用 API 函数clGetPlatformIDs 获取可用 platform 的数量：

cl_int status = 0;cl_uint numPlatforms;cl_platform_id platform = NULL;status = clGetPlatformIDs( 0, NULL, &numPlatforms);if(status != CL_SUCCESS){    printf("Error: Getting Platforms\n");    return EXIT_FAILURE;}

然后根据数量来分配内存，并得到所有可用的 platform，所使用的 API 还是 clGetPlatformIDs。在 OpenCL 中，类似这样的函数调用很常见：第一次调用以取得数目，便于分配足够的内存；然后调用第二次以获取真正的信息。

if (numPlatforms > 0) {    cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));    status = clGetPlatformIDs(numPlatforms, platforms, NULL);    if (status != CL_SUCCESS) {        printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");        return -1;    }

现在，所有的 platform 都存在了变量 platforms 中，接下来需要做的就是取得我们所需的 platform。本人的PC上配置的是 Intel 处理器和 AMD 显卡，专业点的说法叫 Intel 的 CPU 和 AMD 的 GPU :)。所以我这儿有两套 platform，为了体验下 GPU 的快感，所以使用 AMD 的 platform。通过使用clGetPlatformInfo 来获得 platform 的信息。通过这个 API 可以知晓 platform 的厂商信息，以便我们选出需要的 platform。代码如下：

for (unsigned int i = 0; i < numPlatforms; ++i) {        char pbuff[100];        status = clGetPlatformInfo(                     platforms[i],                     CL_PLATFORM_VENDOR,                     sizeof(pbuff),                     pbuff,                     NULL);        platform = platforms[i];        if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {            break;        }    }

不同的厂商信息可以参考 OpenCL Specifications，我这儿只是简单的筛选出 AMD 。

在 platform 上建立 context

第一步是通过 platform 得到相应的 context properties

// 如果我们能找到相应平台，就使用它，否则返回NULLcl_context_properties cps[3] = {    CL_CONTEXT_PLATFORM,    (cl_context_properties)platform,    0};cl_context_properties *cprops = (NULL == platform) ? NULL : cps;

第二步是通过 clCreateContextFromType 函数创建 context。

// 生成 contextcl_context context = clCreateContextFromType(                         cprops,                         CL_DEVICE_TYPE_GPU,                         NULL,                         NULL,                         &status);if (status != CL_SUCCESS) {    printf("Error: Creating Context.(clCreateContexFromType)\n");    return EXIT_FAILURE;} 函数的第二个参数可以设定 context 关联的设备类型。本例使用的是 GPU 作为OpenCL计算设备。目前可以使用的类别包括：- CL_DEVICE_TYPE_CPU- CL_DEVICE_TYPE_GPU- CL_DEVICE_TYPE_ACCELERATOR- CL_DEVICE_TYPE_DEFAULT- CL_DEVICE_TYPE_ALL

在 context 上查询 device

context 创建好之后，要做的就是查询可用的 device。

status = clGetContextInfo(context,                          CL_CONTEXT_DEVICES,                          0,                          NULL,                          &deviceListSize);if (status != CL_SUCCESS) {    printf("Error: Getting Context Info device list size, clGetContextInfo)\n");    return EXIT_FAILURE;}cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);if (devices == 0) {    printf("Error: No devices found.\n");    return EXIT_FAILURE;}status = clGetContextInfo(context,                          CL_CONTEXT_DEVICES,                          deviceListSize,                          devices,                          NULL);if (status != CL_SUCCESS) {    printf("Error: Getting Context Info (device list, clGetContextInfo)\n");    return EXIT_FAILURE;}

与获取 platform 类似，我们调用两次 clGetContextInfo 来完成查询。第一次调用获取关联 context 的 device 个数，并根据个数申请内存；第二次调用获取所有 device 实例。如果你想了解每个 device 的具体信息，可以调用clGetDeviceInfo 函数来获取，返回的信息有设备类型、生产商以及设备对某些扩展功能的支持与否等等。详细使用情况请参阅OpenCL Specifications。

到此，platform 相关的程序已经准备就绪了，下面到此的完整代码：

/* OpenCL_01.cpp  * (c) by keyring <keyrings@163.com> * 2013.10.26 */#if defined(__APPLE__) || defined(__MACOSX)#include <OpenCL/cl.hpp>#else#include <CL/cl.h>#endif#include <iostream>int main(int argc, char const *argv[]){    printf("hello OpenCL\n");    cl_int status = 0;    size_t deviceListSize;    // 得到并选择可用平台    cl_uint numPlatforms;    cl_platform_id platform = NULL;    status = clGetPlatformIDs(0, NULL, &numPlatforms);    if (status != CL_SUCCESS) {        printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");        return EXIT_FAILURE;    }    if (numPlatforms > 0) {        cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));        status = clGetPlatformIDs(numPlatforms, platforms, NULL);        if (status != CL_SUCCESS) {            printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");            return -1;        }// 遍历所有 platform，选择你想用的        for (unsigned int i = 0; i < numPlatforms; ++i) {            char pbuff[100];            status = clGetPlatformInfo(                         platforms[i],                         CL_PLATFORM_VENDOR,                         sizeof(pbuff),                         pbuff,                         NULL);            platform = platforms[i];            if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {                break;            }        }        delete platforms;    }    // 如果我们能找到相应平台，就使用它，否则返回NULL    cl_context_properties cps[3] = {        CL_CONTEXT_PLATFORM,        (cl_context_properties)platform,        0    };    cl_context_properties *cprops = (NULL == platform) ? NULL : cps;    // 生成 context    cl_context context = clCreateContextFromType(                             cprops,                             CL_DEVICE_TYPE_GPU,                             NULL,                             NULL,                             &status);    if (status != CL_SUCCESS) {        printf("Error: Creating Context.(clCreateContexFromType)\n");        return EXIT_FAILURE;    }    // 寻找OpenCL设备    // 首先得到设备列表的长度    status = clGetContextInfo(context,                              CL_CONTEXT_DEVICES,                              0,                              NULL,                              &deviceListSize);    if (status != CL_SUCCESS) {        printf("Error: Getting Context Info device list size, clGetContextInfo)\n");        return EXIT_FAILURE;    }    cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);    if (devices == 0) {        printf("Error: No devices found.\n");        return EXIT_FAILURE;    }    // 然后得到设备列表    status = clGetContextInfo(context,                              CL_CONTEXT_DEVICES,                              deviceListSize,                              devices,                              NULL);    if (status != CL_SUCCESS) {        printf("Error: Getting Context Info (device list, clGetContextInfo)\n");        return EXIT_FAILURE;    }

Running time

前面写了一大篇，其实还没真正进入具体的程序逻辑中，顶多算配好了 OpenCL 运行环境。真正的逻辑代码，即程序的任务就是运行时模块。本例的任务是在一个 4×4的二维空间上，按一定的规则给每个元素赋值，具体代码如下：

#define KERNEL(...)#__VA_ARGS__const char *kernelSourceCode = KERNEL(                                   __kernel void hellocl(__global uint *buffer){    size_t gidx = get_global_id(0);    size_t gidy = get_global_id(1);    size_t lidx = get_local_id(0);    buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);}                               );

这一段就是我们真正的逻辑，也就是代码要干的事。使用的是 OpenCL 自定的一门类C语言，具体的语法什么的现在先不纠结。这段代码是直接嵌入我们的 cpp 文件的静态字符串。你也可以将 kernel 程序单独写成一个文件。

加载 OpenCL 内核程序并创建一个 program 对象

接下来要做的就是读入 OpenCL kernel 程序并创建一个 program 对象。

size_t sourceSize[] = {strlen(kernelSourceCode)};cl_program program = clCreateProgramWithSource(context,                     1,                     &kernelSourceCode,                     sourceSize,                     &status);if (status != CL_SUCCESS) {    printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");    return EXIT_FAILURE;}

本例中的 kernel 程序是作为静态字符串读入的（单独的文本文件也一样），所以使用的是 clCreateProgramWithSource，如果你不想让 kernel 程序让其他人看见，可以先生成二进制文件，再通过clCreateProgramWithBinary 函数动态读入二进制文件，做一定的保密。详细请参阅OpenCL Specifications。

为指定的 device 编译 program 中的 kernel

kernel 程序读入完毕，要做的自然是使用 clBuildProgram 编译 kernel：

status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);if (status != CL_SUCCESS) {    printf("Error: Building Program (clBuildingProgram)\n");    return EXIT_FAILURE;}

最终，kernel 将被相应 device 上的 OpenCL 编译器编译成可执行的机器码。

创建指定名字的 kernel 对象

成功编译后，可以通过 clCreateKernel 来创建一个 kernel 对象。

cl_kernel kernel = clCreateKernel(program, "hellocl", &status);if (status != CL_SUCCESS) {    printf("Error: Creating Kernel from program.(clCreateKernel)\n");    return EXIT_FAILURE;}

引号中的 hellocl 就是 kernel 对象所关联的 kernel 函数的函数名。要注意的是，每个 kernel 对象必须关联且只能关联一个包含于相应 program 对象内的 kernel 程序。实际上，用户可以在 cl 源代码中写任意多个 kernel 程序，但在执行某个 kernel 程序之前必须先建立单独的 kernel 对象，即多次调用clCreateKernel 函数。

为 kernel 创建内存对象

OpenCL 内存对象是指在 host 中创建，用于 kernel 程序的内存类型。按维度可以分为两类，一类是 buffer，一类是 image。buffer 是一维的，image 可以是二维、三维的 texture、frame-buffer 或 image。本例仅仅使用buffer，可以通过 clCreateBuffer 函数来创建。

cl_mem outputBuffer = clCreateBuffer(context, CL_MEM_ALLOC_HOST_PTR, 4 * 4 * 4, NULL, &status);if (status != CL_SUCCESS) {    printf("Error: Create Buffer, outputBuffer. (clCreateBuffer)\n");    return EXIT_FAILURE;}

为 kernel 设置参数

使用 clSetKernelArg 函数为 kernel 设置参数。传递的参数既可以是常数，变量，也可以是内存对象。本例传递的就是内存对象。

status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&outputBuffer);if (status != CL_SUCCESS) {    printf("Error: Setting kernel argument. (clSetKernelArg)\n");    return EXIT_FAILURE;}

该函数每次只能设置一个参数，如有多个参数，需多次调用。而且 kernel 程序中所有的参数都必须被设置，否则在启动 kernel 程序是会报错。指定位置的参数的类型最好和对应 kernel 函数内参数类型一致，以免产生各种未知的错误。在设置好指定参数后，每次运行该 kernel 程序都会使用设置值，直到用户使用次 API 重新设置参数。

在指定的 device 上创建 command queue

command queue 用于光里将要执行的各种命令。可以通过 clCreateCommandQueue 函数创建。其中的 device 必须为 context 的关联设备，所有该 command queue 中的命令都会在这个指定的 device 上运行。

cl_command_queue commandQueue = clCreateCommandQueue(context,                                devices[0],                                0,                                &status);if (status != CL_SUCCESS) {    printf("Error: Create Command Queue. (clCreateCommandQueue)\n");    return EXIT_FAILURE;}

将要执行的 kernel 放入 command queue

创建好 command queue 后，用户可以创建相应的命令并放入 command queue 中执行。OpenCL 提供了三种方案来创建 kernel 执行命令。最常用的即为本例所示的运行在指定工作空间上的 kernel 程序，使用了clEnqueueNDRangeKernel 函数。

size_t globalThreads[] = {4, 4};size_t localThreads[] = {2, 2};status = clEnqueueNDRangeKernel(commandQueue, kernel,                                2, NULL, globalThreads,                                localThreads, 0,                                NULL, NULL);if (status != CL_SUCCESS) {    printf("Error: Enqueueing kernel\n");    return EXIT_FAILURE;}

clEnqueueNDRangeKernel 函数每次只能将一个 kernel 对象放入 command queue 中，用户可以多次调用该 API 将多个 kernel 对象放置到一个 command queue 中，command queue 中的不同 kernel 对象的工作区域完全不相关。其余两个 APIclEnqueueTask 和 clEnqueueNativeKernel 的用法就不多讲了，详情请参阅 OpenCL Specificarions。

最后可以用 clFinish 函数来确认一个 command queue 中所有的命令都执行完毕。函数会在 command queue 中所有 kernel 执行完毕后返回。

// 确认 command queue 中所有命令都执行完毕status = clFinish(commandQueue);if (status != CL_SUCCESS) {    printf("Error: Finish command queue\n");    return EXIT_FAILURE;}

将结果读回 host

计算完毕，将结果读回 host 端。使用 clEnqueueReadBuffer 函数将 OpenCL buffer 对象中的内容读取到 host 可以访问的内存空间。

// 将内存对象中的结果读回Hoststatus = clEnqueueReadBuffer(commandQueue,                             outputBuffer, CL_TRUE, 0,                             4 * 4 * 4, outbuffer, 0, NULL, NULL);if (status != CL_SUCCESS) {    printf("Error: Read buffer queue\n");    return EXIT_FAILURE;}

当然，为了看下程序的运行效果，咱们当然得看看运行结果啦。打印一下吧：

// Host端打印结果printf("out:\n");for (int i = 0; i < 16; ++i) {    printf("%x ", outbuffer[i]);    if ((i + 1) % 4 == 0)        printf("\n");}

资源回收

程序的最后是对所有创建的对象进行释放回收，与C/C++的内存回收同理。

// 资源回收status = clReleaseKernel(kernel);status = clReleaseProgram(program);status = clReleaseMemObject(outputBuffer);status = clReleaseCommandQueue(commandQueue);status = clReleaseContext(context);free(devices);delete outbuffer;

总结

这次使用一个小例子来详细说明了 OpenCL 编程的一般步骤。其实这些步骤一般都是固定的。真正需要我们注意的是 OpenCL Kernel 程序的编写。当然，合理高效的利用 API 也是一门技术活。

最后给出本实例的全部代码：

/*OpenCL_01.cpp *(c) by keyring <keyrings@163.com> *2013.10.26 */#include <iostream>#if defined(__APPLE__) || defined(__MACOSX)#include <OpenCL/cl.hpp>#else#include <CL/cl.h>#endif#define KERNEL(...)#__VA_ARGS__const char *kernelSourceCode = KERNEL(                                   __kernel void hellocl(__global uint *buffer){    size_t gidx = get_global_id(0);    size_t gidy = get_global_id(1);    size_t lidx = get_local_id(0);    buffer[gidx + 4 * gidy] = (1 << gidx) | (0x10 << gidy);}                               );int main(int argc, char const *argv[]){    printf("hello OpenCL\n");    cl_int status = 0;    size_t deviceListSize;    // 得到并选择可用平台    cl_uint numPlatforms;    cl_platform_id platform = NULL;    status = clGetPlatformIDs(0, NULL, &numPlatforms);    if (status != CL_SUCCESS) {        printf("ERROR: Getting Platforms.(clGetPlatformIDs)\n");        return EXIT_FAILURE;    }    if (numPlatforms > 0) {        cl_platform_id *platforms = (cl_platform_id *)malloc(numPlatforms * sizeof(cl_platform_id));        status = clGetPlatformIDs(numPlatforms, platforms, NULL);        if (status != CL_SUCCESS) {            printf("Error: Getting Platform Ids.(clGetPlatformIDs)\n");            return -1;        }        for (unsigned int i = 0; i < numPlatforms; ++i) {            char pbuff[100];            status = clGetPlatformInfo(                         platforms[i],                         CL_PLATFORM_VENDOR,                         sizeof(pbuff),                         pbuff,                         NULL);            platform = platforms[i];            if (!strcmp(pbuff, "Advanced Micro Devices, Inc.")) {                break;            }        }        delete platforms;    }    // 如果我们能找到相应平台，就使用它，否则返回NULL    cl_context_properties cps[3] = {        CL_CONTEXT_PLATFORM,        (cl_context_properties)platform,        0    };    cl_context_properties *cprops = (NULL == platform) ? NULL : cps;    // 生成 context    cl_context context = clCreateContextFromType(                             cprops,                             CL_DEVICE_TYPE_GPU,                             NULL,                             NULL,                             &status);    if (status != CL_SUCCESS) {        printf("Error: Creating Context.(clCreateContexFromType)\n");        return EXIT_FAILURE;    }    // 寻找OpenCL设备    // 首先得到设备列表的长度    status = clGetContextInfo(context,                              CL_CONTEXT_DEVICES,                              0,                              NULL,                              &deviceListSize);    if (status != CL_SUCCESS) {        printf("Error: Getting Context Info device list size, clGetContextInfo)\n");        return EXIT_FAILURE;    }    cl_device_id *devices = (cl_device_id *)malloc(deviceListSize);    if (devices == 0) {        printf("Error: No devices found.\n");        return EXIT_FAILURE;    }    // 现在得到设备列表    status = clGetContextInfo(context,                              CL_CONTEXT_DEVICES,                              deviceListSize,                              devices,                              NULL);    if (status != CL_SUCCESS) {        printf("Error: Getting Context Info (device list, clGetContextInfo)\n");        return EXIT_FAILURE;    }    // 装载内核程序，编译CL program ,生成CL内核实例    size_t sourceSize[] = {strlen(kernelSourceCode)};    cl_program program = clCreateProgramWithSource(context,                         1,                         &kernelSourceCode,                         sourceSize,                         &status);    if (status != CL_SUCCESS) {        printf("Error: Loading Binary into cl_program (clCreateProgramWithBinary)\n");        return EXIT_FAILURE;    }    // 为指定的设备编译CL program.    status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);    if (status != CL_SUCCESS) {        printf("Error: Building Program (clBuildingProgram)\n");        return EXIT_FAILURE;    }    // 得到指定名字的内核实例的句柄    cl_kernel kernel = clCreateKernel(program, "hellocl", &status);    if (status != CL_SUCCESS) {        printf("Error: Creating Kernel from program.(clCreateKernel)\n");        return EXIT_FAILURE;    }    // 创建 OpenCL buffer 对象    unsigned int *outbuffer = new unsigned int [4 * 4];    memset(outbuffer, 0, 4 * 4 * 4);    cl_mem outputBuffer = clCreateBuffer(    context,     CL_MEM_ALLOC_HOST_PTR,     4 * 4 * 4,     NULL,     &status);    if (status != CL_SUCCESS) {        printf("Error: Create Buffer, outputBuffer. (clCreateBuffer)\n");        return EXIT_FAILURE;    }    //  为内核程序设置参数    status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&outputBuffer);    if (status != CL_SUCCESS) {        printf("Error: Setting kernel argument. (clSetKernelArg)\n");        return EXIT_FAILURE;    }    // 创建一个OpenCL command queue    cl_command_queue commandQueue = clCreateCommandQueue(context,                                    devices[0],                                    0,                                    &status);    if (status != CL_SUCCESS) {        printf("Error: Create Command Queue. (clCreateCommandQueue)\n");        return EXIT_FAILURE;    }    // 将一个kernel 放入 command queue    size_t globalThreads[] = {4, 4};    size_t localThreads[] = {2, 2};    status = clEnqueueNDRangeKernel(commandQueue, kernel,                                    2, NULL, globalThreads,                                    localThreads, 0,                                    NULL, NULL);    if (status != CL_SUCCESS) {        printf("Error: Enqueueing kernel\n");        return EXIT_FAILURE;    }    // 确认 command queue 中所有命令都执行完毕    status = clFinish(commandQueue);    if (status != CL_SUCCESS) {        printf("Error: Finish command queue\n");        return EXIT_FAILURE;    }    // 将内存对象中的结果读回Host    status = clEnqueueReadBuffer(commandQueue,                                 outputBuffer, CL_TRUE, 0,                                 4 * 4 * 4, outbuffer, 0, NULL, NULL);    if (status != CL_SUCCESS) {        printf("Error: Read buffer queue\n");        return EXIT_FAILURE;    }    // Host端打印结果    printf("out:\n");    for (int i = 0; i < 16; ++i) {        printf("%x ", outbuffer[i]);        if ((i + 1) % 4 == 0)            printf("\n");    }    // 资源回收    status = clReleaseKernel(kernel);    status = clReleaseProgram(program);    status = clReleaseMemObject(outputBuffer);    status = clReleaseCommandQueue(commandQueue);    status = clReleaseContext(context);    free(devices);    delete outbuffer;    system("pause");    return 0;}

0 0