第二章 HelloWorld:一个OpenCL例子

本章将介绍一下概念：

• 选择OpenCL平台并创建一个上下文。
• 列举设备并创建一个命令队列。
• 创建和构建一个程序对象。
• 创建一个内核对象，并为内核参数创建内存对象。
• 执行内核并读取结果。
• 检查OpenCL中的错误。

2.2.1 选择OpenCL平台并创建一个上下文

创建OpenCL程序的第一步是选择一个平台。OpenCL使用可安装客户驱动程序(Installable Client Driver, ICD)模型，一个系统上可以有多个OpenCL实现并存。

//选择平台并创建上下文

cl_context CreateContext() {
    cl_int errNum;
    cl_uint numPlatforms;
    cl_platform_id firstPlatformId;
    cl_context context = NULL;
　　 errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatfor,s);
　　 if (err != CL_SUCCESS || numPlatforms <= 0) {
　　　　　　cerr << "Failed to find any OpenCL platforms." << endl;
　　　　　　return NULL;
　　 }
　　 //在这个平台上创建上下文，首先尝试创建一个基于GPU的上下文，如果失败了，尝试创建一个基于CPU的上下文
　　 cl_context_properties contextProperties[] = { 
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　CL_CONTEXT_PLATFORM,
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　(cl_context_properties)firstPlatformId,
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　0
　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　}
　　　context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
　　  if (errNum != CL_SUCCESS) {
　　　      cout << "Could not create GPU context, trying CPU...." << endl;
　　　　　　 context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
　　　　　　 if (errNum != CL_SUCCESS) {
　　　　　　　　　　cerr << "Failed to create an OpenCL CPU or CPU context." << endl;
                 return NULL;
　　　　　　 }
      }
      return context;
}

2.2.2 选择设备并创建命令队列　　

选择平台并创建一个上下文之后，HelloWorld应用的下一步要选择一个设备，并创建一个命令队列。设备在计算机硬件底层，如CPU或GPU。要与设备通信，应用程序必须为它创建一个命令队列。将在设备上完成的操作要在命令队列中排队。

//选择第一个可用设备并创建一个命令队列
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 (errNum != CL_SUCESS) {
　　　　cerr << "Failed ..." << endl;
　　　　return NULL;
　　}
　　if (deviceBufferSize <= 0) {
　　　　cerr << "No devices available." << endl;
　　　　return NULL;
　　}
　　devices = new cl_device_id[deviceBufferSize/sizeof(cl_device_id)];
　　errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL):
　　if (errNum != CL_SUCCESS) {
　　　　cerr << "Failed to get device IDs" << endl;
　　　　return NULL;
　　}
　　//我们先只选择第一个可用的设备
　　commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
　　if (commandQueue == NULL) {
　　　　cerr << "Failed to create commandQueue for device 0" << endl;
　　　　return NULL;
　　}
　　*device = devices[0];
　　delete [] devices;
　　return commandQueue;
}

2.2.3 创建和构建程序对象　

下一步是从HelloWorld.cl文件加载OpenCL C内核源代码，由它创建一个程序对象。这个程序对象用内核源代码加载，然后进行编译，从而在与上下文的关联的设备上执行。一般地，OpenCL中的程序对象会存储于上下文相关联的所有设备的已编译可执行代码。

//从磁盘加载内核源文件并创建和构建一个程序对象
cl_program CreateProgram(cl_context context, cl_device_id device, const char* filename) {
    cl_int errNum;
    cl_program program;
    ifstream kernelFile(flieName, ios::in);        
    if (!kernelFile.is_open()) {
        cerr << "Failed to open file for reading:" << filename << endl;
        return NULL;    
    }
    ostringstream oss;
    oss << kernelFile.rebuf();
    string srcStdStr = oss.str();
    const char* srcStr = srcStdStr.c_str();
    program = clCreateProgramWithSource(context, 1, (const char**)&srcStr, NULL, NULL);
    if (program == NULL) {
        cerr << "Failed to create CL program from source" << endl;
        return NULL;
    }
    errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    if (errNum != CL_SUCCESS) {
        char buildLog[16384];
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, sizeof(buildLog), buildLog, NULL);
        cerr << "Error in kernel:" << endl;
        cerr << buildLog;
        clReleaseProgram(program);
        return NULL;
    }
    return program;
}　　

 从磁盘加载.cl文件，并存储在一个字符串中。然后通过调用clCreateProgramWithSource()创建程序对象，它会由内核源代码创建程序对象。创建了程序对象之后，通过调用clBuildProgram()编译内核源代码。这个函数会为关联的设备编译内核，如果编译成功，则把编译代码存储在程序对象中。如果编译过程失败，可以使用clGetProgramBuildInfo()获取构建日志。构建日志中的字符串包含了OpenCL内核编译过程中生成的所有编译错误。

2.2.4 创建内核和内存对象

要执行OpenCL计算内核，需要再内存中分配内核函数的参数，以便在OpenCL设备上访问。本例子的内核是一个简单的函数，它要计算两个数组(a和b）中各个元素之和，并把结果存储在另一个数组(result)中。创建一个内核对象，将其编译到程序对象中。另外，将分配数组(a,b和result）并填入数据。在宿主机内存中创建这些数组之后，调用CreateMemObjects()，它会把这些数组复制到内存对象，然后传入内核。

//创建内核
//Create OpenCL kernel
kernel = clCreateKernel(program, "hello_kernel", NULL);
if (kernel == NULL) {
    cerr << "Failed to create kernel" << endl;
    Cleanup(context, commandQueue, program, kernel, memObjects);
    return 1;
}
float resutl[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)(i * 2);
}
if (!CreateMemObjects(context, memObjects, a, b) {
    Cleanup(context, commandQueue, program, kernel, memObjects);
    return 1;
}

CreateMemObjects()函数为各个数组分别调用clCreateBuffer()来创建一个内存对象。内核对象分配在设备内存中，可以由内核函数直接访问。对于输入数组(a和b），缓冲区CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR内存类型来创建，这说明这些数组对内核是只读的，可以从宿主机内存复制到设备内存。数组本身作为参数传递到clCreateBuffer(),这会将数组的内容复制到设备上为内存对象分配的存储空间中。result数组用内存类型CL_MEM_READ_WRITE创建，这说明这个数组对内核是可读、可写的。

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);
    if (memObjects[0] == NULL || memObjects[1] == NULL || memObjects[2] == NULL) {
        //to do
    }
    return true;  
}

2.2.5 执行内核

内核函数的所有参数需要使用clSetKernelArg()设置。这个函数的第一个参数时参数的索引。hello_kernel()有3个参数(a,b,result),分别对应索引0,1,2。CreateMemObjects()创建的内存对象传入内核对象

//为内核函数设置参数
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);

errNum = clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE, 0, ARRAY_SIZE * sizeof(float), result, 0 , NULL, NULL);