OpenCL框架

OpenCL是一个异构平台编程框架，也是异构并行计算的行业标准，由Khronos Group来维护。旨在充分利用GPU强大的并行计算能力以及CPU的协同工作，更高效的利用硬件完成大规模(尤其是并行度高的)计算。

简单说，CPU并不适合计算，它是多指令单数据流(MISD)的体系结构，更加擅长的是做逻辑控制，而数据处理基本是单流水线的，所以像代码for(i=0;...;i++)这种在CPU上要重复迭代跑很多遍。但是我们的显卡GPU则不这样，GPU是典型的单指令多数据(SIMD)的体系结构，它不擅长逻辑控制，但是确实天生的向量计算机器，对于上面for(i=0;...;i++)这样的代码有时只需要跑一遍，所以图形世界中那么多的顶点、片段才能快速的并行在显卡中渲染处理。

• OpenCL架构

代码由两部分构成：

1. Host: 主程序，main入口代码，在CPU上运行。
2. Devices: 各个设备(GPU)上运行的kernel程序，类似opengl的shader程序。

很像OpenGL的渲染管线，在主程序创建上下文环境，分配共享内存(CPU/GPU之间交互数据)，加载cl程序，发送到GPU设备上执行，然后主程序取出结果数据，大体就这么个流程。

OpenCL编译环境配置(VS+Intel)

1. 安装对应显卡的opencl的sdk，我测试的是intel显卡(下载)，一路点下一步安装没啥特别的，记住安装位置。

   

2. 在vs创建空项目，然后属性页包含sdk头文件和链接OpenCL.lib库。

   

   

   

示例程序

下面是两个数组相加的示例程序。

• 主程序

#include <CL/cl.h>
#include <fstream>
#include <iostream>
#include <sstream>

const int ARRAY_SIZE = 1000;

// 一、 选择OpenCL平台并创建一个上下文
cl_context CreateContext()
{
    cl_int errNum;
    cl_uint numPlatforms;
    cl_platform_id firstPlatformId;
    cl_context context = NULL;

    // 选择可用的平台中的第一个
    errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
    if (errNum != CL_SUCCESS || numPlatforms <= 0) {
        std::cerr << "Failed to find any OpenCL platforms." << std::endl;
        return NULL;
    }

    // 创建一个OpenCL上下文环境
    cl_context_properties contextProperties[] = {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)firstPlatformId,
        0
    };
    context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
        NULL, NULL, &errNum);

    return context;
}

// 二、 创建设备并创建命令队列
cl_command_queue CreateCommandQueue(cl_context context, cl_device_id* device)
{
    cl_int errNum;
    cl_device_id* devices;
    cl_command_queue commandQueue = NULL;
    size_t deviceBufferSize = -1;

    // 获取设备缓冲区大小
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);

    if (deviceBufferSize <= 0) {
        std::cerr << "No devices available.";
        return NULL;
    }

    // 为设备分配缓存空间
    devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);

    // 选取可用设备中的第一个
    commandQueue = clCreateCommandQueueWithProperties(context, devices[0], 0, NULL);

    *device = devices[0];
    delete[] devices;
    return commandQueue;
}

// 三、创建和构建程序对象
cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
    cl_int errNum;
    cl_program program;

    std::ifstream kernelFile(fileName, std::ios::in);
    if (!kernelFile.is_open()) {
        std::cerr << "Failed to open file for reading: " << fileName << std::endl;
        return NULL;
    }

    std::ostringstream oss;
    oss << kernelFile.rdbuf();

    std::string srcStdStr = oss.str();
    const char* srcStr = srcStdStr.c_str();
    program = clCreateProgramWithSource(context, 1,
        (const char**)&srcStr,
        NULL, NULL);

    errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

    return program;
}

// 创建和构建程序对象
bool CreateMemObjects(cl_context context, cl_mem memObjects[3],
    float* a, float* b)
{
    memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        sizeof(float) * ARRAY_SIZE, a, NULL);
    memObjects[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        sizeof(float) * ARRAY_SIZE, b, NULL);
    memObjects[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
        sizeof(float) * ARRAY_SIZE, NULL, NULL);
    return true;
}

// 释放OpenCL资源
void Cleanup(cl_context context, cl_command_queue commandQueue,
    cl_program program, cl_kernel kernel, cl_mem memObjects[3])
{
    for (int i = 0; i < 3; i++) {
        if (memObjects[i] != 0)
            clReleaseMemObject(memObjects[i]);
    }
    if (commandQueue != 0)
        clReleaseCommandQueue(commandQueue);

    if (kernel != 0)
        clReleaseKernel(kernel);

    if (program != 0)
        clReleaseProgram(program);

    if (context != 0)
        clReleaseContext(context);
}

int main(int argc, char** argv)
{
    cl_context context = 0;
    cl_command_queue commandQueue = 0;
    cl_program program = 0;
    cl_device_id device = 0;
    cl_kernel kernel = 0;
    cl_mem memObjects[3] = { 0, 0, 0 };
    cl_int errNum;

    // 一、选择OpenCL平台并创建一个上下文
    context = CreateContext();

    // 二、 创建设备并创建命令队列
    commandQueue = CreateCommandQueue(context, &device);

    // 三、创建和构建程序对象
    program = CreateProgram(context, device, "HelloWorld.cl");

    // 四、 创建OpenCL内核并分配内存空间
    kernel = clCreateKernel(program, "hello_kernel", NULL);

    // 创建要处理的数据
    float result[ARRAY_SIZE];
    float a[ARRAY_SIZE];
    float b[ARRAY_SIZE];
    for (int i = 0; i < ARRAY_SIZE; i++) {
        a[i] = (float)i;
        b[i] = (float)(ARRAY_SIZE - i);
    }

    // 创建内存对象
    if (!CreateMemObjects(context, memObjects, a, b)) {
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

    // 五、 设置内核数据并执行内核
    errNum = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memObjects[0]);
    errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObjects[1]);
    errNum |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &memObjects[2]);

    size_t globalWorkSize[1] = { ARRAY_SIZE };
    size_t localWorkSize[1] = { 1 };

    errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL,
        globalWorkSize, localWorkSize,
        0, NULL, NULL);

    // 六、 读取执行结果并释放OpenCL资源
    errNum = clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE,
        0, ARRAY_SIZE * sizeof(float), result,
        0, NULL, NULL);

    for (int i = 0; i < ARRAY_SIZE; i++) {
        std::cout << result[i] << " ";
    }
    std::cout << std::endl;
    std::cout << "Executed program succesfully." << std::endl;
    getchar();
    Cleanup(context, commandQueue, program, kernel, memObjects);

    return 0;
}

• cl程序(HelloWorld.cl和主程序源码放一起)

__kernel void hello_kernel(__global const float *a,
	__global const float *b,
	__global float *result)
{
	int gid = get_global_id(0);
 
	result[gid] = a[gid] + b[gid];
}