OpenCL向量相加

原文http://www.olcf.ornl.gov/training_articles/opencl-vector-addition/

本文仅仅是为了学习OpenCL而做的的相关翻译。

由于原文中的例子不能在我的环境中运行，因此做了一些改动。

通过这个例子能很好地了解OpenCL的编程模型。

1. 简介

这个例子是表示了两个向量相加，可以认为是OpenCL中的"hello world"。为了使程序更容易理解，没有加入错误处理机制。

//vecAdd.c

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <CL/opencl.h>

 
// OpenCL kernel. Each work item takes care of one element of c
const char *kernelSource =                                      "\n" \
"__kernel void vecAdd(  __global float *a,                       \n" \
"                       __global float *b,                       \n" \
"                       __global float *c,                       \n" \
"                       const unsigned int n)                    \n" \
"{                                                               \n" \
"    //Get our global thread ID                                  \n" \
"    int id = get_global_id(0);                                  \n" \
"                                                                \n" \
"    //Make sure we do not go out of bounds                      \n" \
"    if (id < n)                                                 \n" \
"        c[id] = a[id] + b[id];                                  \n" \
"}                                                               \n" \
                                                                "\n" ;
 
int main( int argc, char* argv[] )
{
    // 向量长度
    int n = 8;
 
    // 输入向量
    int *h_a;
    int *h_b;
    // 输出向量
    int *h_c;
 
    // 设备输入缓冲区
    cl_mem d_a;
    cl_mem d_b;
    // 设备输出缓冲区
    cl_mem d_c;
 
    cl_platform_id cpPlatform;        // OpenCL 平台
    cl_device_id device_id;           // device ID
    cl_context context;               // context
    cl_command_queue queue;           // command queue
    cl_program program;               // program
    cl_kernel kernel;                 // kernel
 
    //（每个向量的字节数）
    size_t bytes = n*sizeof(int);
 
    //（为每个向量分配内存）
    h_a = (int*)malloc(bytes);
    h_b = (int*)malloc(bytes);
    h_c = (int*)malloc(bytes);
 
    //（初始化向量）
	int i;
    for( i = 0; i < n; i++ )
    {
        h_a[i] = i;
        h_b[i] = i;
    }

	size_t globalSize, localSize;
    cl_int err;
 
    //（每个工作组的工作节点数目）
    localSize = 2;
 
    //（所有的工作节点）
    globalSize = (size_t)ceil(n/(float)localSize)*localSize;
	printf("%d\n",globalSize);
 
    //（获得平台ID）
    err = clGetPlatformIDs(1, &cpPlatform, NULL);
    //（获得设备ID，与平台有关）
    err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
 
    //（根据设备ID，得到上下文）
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
 
    //（根据上下文，在设备上创建命令队列）
    queue = clCreateCommandQueue(context, device_id, 0, &err);
 
    //（根据OpenCL源程序创建计算程序）
    program = clCreateProgramWithSource(context, 1,
                            (const char **) & kernelSource, NULL, &err);
 
    //（创建可执行程序）
    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	
    //（在上面创建的程序中创建内核程序）
    kernel = clCreateKernel(program, "vecAdd", &err);
 
    //（分配设备缓冲）
	d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
    d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
    d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);
 
    // （将向量信息写入设备缓冲）
    err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,
                                   bytes, h_a, 0, NULL, NULL);
    err |= clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,
                                   bytes, h_b, 0, NULL, NULL);
 
    // （设置计算内核的参数）
    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);
    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);
    err |= clSetKernelArg(kernel, 3, sizeof(int), &n);
 
    // （在数据集的范围内执行内核）Execute the kernel over the entire range of the data set
    err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,
                                                              0, NULL, NULL);
 
    // （在读出结果之前，等待命令队列执行完毕）Wait for the command queue to get serviced before reading back results
    clFinish(queue);
 
    // （从设备缓冲区读出结果）Read the results from the device
    clEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,
                                bytes, h_c, 0, NULL, NULL );
 
    //（输出读出的结果）
    float sum = 0;
    for(i=0; i<n; i++)
		printf("%d ",h_c[i]);
 
    // （释放资源）
    clReleaseMemObject(d_a);
    clReleaseMemObject(d_b);
    clReleaseMemObject(d_c);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(queue);
    clReleaseContext(context);
 
    //（释放内存）
    free(h_a);
    free(h_b);
    free(h_c);
	system("pause");
    return 0;
}

2. 基本解释

2.1 内核：

Kernel是OpenCL代码的核心。全部kernel必须要作为一个C字符串读入，最容易的方式就是将整个kernel用引号包起来，行尾回车。在真正的程序中，应该将kernel放在一个独立的文件中。

// OpenCL kernel. Each work item takes care of one element of c
const char *kernelSource =                                      "\n" \
"__kernel void vecAdd(  __global float *a,                       \n" \
"                       __global float *b,                       \n" \
"                       __global float *c,                       \n" \
"                       const unsigned int n)                    \n" \
"{                                                               \n" \
"    //Get our global thread ID                                  \n" \
"    int id = get_global_id(0);                                  \n" \
"                                                                \n" \
"    //Make sure we do not go out of bounds                      \n" \
"    if (id < n)                                                 \n" \
"        c[id] = a[id] + b[id];                                  \n" \
"}                                                               \n" \
                                                                "\n" ;

　　

查看一下这个简单的内核由什么内容组成：

__kernel void vecAdd(  __global float *a, __global float *b,
                       __global float *c, const unsigned int n)

__kernel 指明这是一个OpenCL内核，__global 说明指针指向的是全局的设备内存空间，其它的就是C语言的函数的语法。kernel必须返回空类型。

int id = get_global_id(0);

得到第0维全局工作节点的ID。

if (id < n)
    c[id] = a[id] + b[id];

工作组的数目必须是一个整数，或者每个工作组的工作节点数目必须能被全部工作节点数目整除。由于共组的的大小被用来协调性能，没有必要一定能被所有线程数目整除，所以通常启用的线程比所需要的线程多一些，并忽略掉多余的。在考察了问题域之后，就能访问、操作设备内存了。
2.2 内存：

    // 输入向量
    int *h_a;
    int *h_b;
    // 输出向量
    int *h_c;
 
    // 设备输入缓冲区
    cl_mem d_a;
    cl_mem d_b;
    // 设备输出缓冲区
    cl_mem d_c;

CPU和GPU有不同的内存空间，所以必须支持对内存分别引用，一个集市主机数组指针，另外一个集是设备内存的操作句柄。这儿我们用 h_和d_前缀来区分。

2.3 线程映射：

    //（每个工作组的工作节点数目）
    localSize = 2;
 
    //（所有的工作节点）
    globalSize = (size_t)ceil(n/(float)localSize)*localSize;

为了将问题映射到底层硬件，必须指明局部的大小，和全局的大小。局部大小定义了工作组中节点的数目，子NVIDIA GPU上这相当于线程块内线程的数目。全局大小定义了所有启动的工作节点数目。localSize大小必须能被globalSize整除，所以我们计算了一个最小的整数能覆盖问题域，并且能被localSize整除。

2.4 环境配置：

    //（绑定平台）
    err = clGetPlatformIDs(1, &cpPlatform, NULL);

每个硬件提供商都有不同的平台，在用之前就应该给定，这儿clGetPlatformIDs()会将cpPlatform赋予系统可用的平台。例如，如果系统包含了AMD CPU和NVIDIA GPU，且这两个平台都安装了合适的OpenCL驱动，那这里平台都是可用的。（注：要使用不同的平台驱动，必须安装相关的驱动，在本例中我安装了AMD(ATI)的app SDK v2.5和Intel的intel_ocl_sdk_1.5_runtime_setup，所以会有两个平台，但是由于我的ATI的显卡GPU不能被app SDK v2.5支持，所以的获得设备ID时没有用GPU设备，而是用了CPU设备。如果这里配置不正确，下面的可能就无法进行）

//（获得设备ID，与平台有关）
err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);

可以查询平台来得到它包含什么样的设备。在这个例子中，用枚举值CL_DEVICE_TYPE_CPU来查询平台上的CPU设备。

    //（根据上下文，在设备上创建命令队列）
    queue = clCreateCommandQueue(context, device_id, 0, &err);

在使用OpenCL设备之前，必须要配置context，context被用来管理命令队列，内存和内核的活动。一个context可以包含不止一个设备。

    //（根据上下文，在设备上创建命令队列）
    queue = clCreateCommandQueue(context, device_id, 0, &err);

命令队列被用来将命令从主机放入指定的设备。内存的转移和内核的活动都能被放入命令队列在合适的时候在指定的设备上执行。

2.5 编译内核：

    //（根据OpenCL源程序创建计算程序）
    program = clCreateProgramWithSource(context, 1,
                            (const char **) & kernelSource, NULL, &err);
 
    //（创建可执行程序）
    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
    
    //（在上面创建的程序中创建内核程序）
    kernel = clCreateKernel(program, "vecAdd", &err);

为了保证对于大多数设备的可移植性，默认运行内核的方式就是用即时（Just-in-time）编译我们必须为给定上下文的设备准备源码。首先，创建程序，这是一个内核程序的的集合，然后根据程序来创建各自的内核程序。

2.6 准备数据：

    //（分配设备缓冲）
    d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
    d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
    d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);
 
    // （将向量信息写入设备缓冲）
    err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,
                                   bytes, h_a, 0, NULL, NULL);
    err |= clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,
                                   bytes, h_b, 0, NULL, NULL);
 
    // （设置计算内核的参数）
    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);
    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);
    err |= clSetKernelArg(kernel, 3, sizeof(int), &n);

在启动内核之前，必须在设备和主机之间创建缓冲区，将主机数据绑定到新创建的设备缓冲区上，最后，设置内核参数。

2.7 启动内核：

// （在数据集的范围内执行内核）
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,
0, NULL, NULL);

一旦内存驻留在设备上之后，内核就能排队启动了。

2.8 取回结果：

    // （在读出结果之前，等待命令队列执行完毕）
    clFinish(queue);
 
    // （从设备缓冲区读出结果）
    clEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,
                                bytes, h_c, 0, NULL, NULL );

可以进行阻断，直到所有的命令队列执行完毕，然后将设备上的结果取回到主机。

3. 运行环境

3.1  OpenCL:

　　AMD app sdk v2.5

　　intel_ocl_sdk_1.5_runtime

3.2 Visual Studio 2010 express