GPGPU.ORG.CN

1、计算设备(Compute Device). 是指系统中支持OpenCL的各类处理器，可以是CPU, GPU，甚至是DSP。一个系统中可以有多个Device连接到主机(Host)上。另外，每个device中会包含多个计算单元(Compute Unit)，例如多核的CPU或者GPU中的多处理器(Multi-Processor). 而每个计算单元中还可以包含多个处理单位(PE, Processing Element). CPU中每个计算单元只含有一个PE, 而GPU中则含有多个。我们可以通过clGetDeviceIDs()这个函数来检查系统中是否有支持OpenCL的设备，并且配合clGetDeviceInfo()函数来了解该设备中如设备类型、计算单元个数等信息。

2、执行(Execution)
最终在Device上执行的代码成为内核(Kernel)，这类代码是C语言的扩展，像下面这样子:

__kernel sum(__global const float* a, __global const float* b,
    __global float* answer)
{
   int xid = get_global_id(0);
    answer[xid] = a[xid] + b[xid];
}

这段代码实现对两个向量a和b求和。__kernel 标识该函数是内核程序，将在device中执行; __global 标识指针所指向内存区域的类型为全局内存(将在存储对象中介绍)。xid则表示向量中的位置. 对于长度为M的向量来说，将会有M个kernel的实例(instance)被创建，每一个实例被称为一个工作单元(Work-Item)，这些工作单元按照用户的设定组成工作组(Work-Group). 我们为这N个单元创建一个索引，则其长度Gx = M，每个工作组的长度为Sx。另外索引的总长度(global size)也称为NDRange(N Dimensional Range)。如果你熟悉CUDA会发现work-item就是thread，work-group就是block:)

  既然叫N Dimensional，N可以为1，2，3. N=2时的情形是这样的：

每个work-item在整个NDRange中的位置是可知，并且唯一的。在kernel中可以通过get_global_id(d)函数来获取该位置，其中d表示维度。示例中d=0表示在第一个维度中的索引位置xid(xid = 0,1,…Sx-1)。
另外OpenCL也提供了get_local_id()函数来获取work-item在work-group中的位置。

3、存储对象(Memory Object)
在OpenCL的模型中，共有四种存储类型：全局内存(Global Memory, __global)，带有缓存的全局内存/常量内存(Constant Memory, __constant)，局部内存(Local Memory, __local)和私有内存(Private Memory, __private). 其中全局内存可供所有的work-item读写，一般全局内存读写速度较慢。而Global/Constant内存由于具有高速缓存(Data Cache)，一般具有更快的读写速度，但对于某一类Device（比如GPU），这类内存通常是只读的Constant Memory（比如纹理内存，Texture）。Local Memory由于其物理位置在处理器芯片上，具有极高的读写速度，但同样的限制了它只能被同一个work-group的work-items共享，并且容量有限。Private Memory是work-item的私有物品，别人看不到也没法用。容量最少，但有最快的访问速度。
 
可以调用如下函数进行内存操作：

// Allocate the OpenCL buffer memory objects for source and result on the device
a = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float)*M, NULL, &err);
answer = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float)*M, NULL, &err);
 
// Asynchronous write of data to GPU device
err = clEnqueueWriteBuffer(cqCommandQue, a, CL_FALSE, 0, sizeof(cl_float)*M, srcA, 0, NULL, NULL);
 
// launch kernel
// clEnqueueNDRangeKernel()
// Synchronous/blocking read of results to host memory->dst
err = clEnqueueReadBuffer(cqCommandQue, answer, CL_TRUE, 0, sizeof(cl_float)*M, dst, 0, NULL, NULL);

在内核中，可以像C语言中一样，使用指针来访问内存，正如我们在kernel示例中看到的一样。需要注意的是，对于GPU的全局内存而言，从host访问是有缓存的，而从kernel访问是没有缓存的。

上一节中我们说到，OpenCL中有一类特殊的内存对象：图像对象(Image Object)
OpenCL中支持的图像可以为1、2、3维，被优化存储在非线性内存中，以便于快速访问。其中1、2维图像由clCreateImage2D()函数创建，而3维图像由clCreateImage3D()函数创建。在创建时，需要说明该图像的格式（RGBA,R,A…）和数据类型（float, int8, unsigned int16…）。
而在kernel中读取图像中的数据，则需要定义一个采样器(sampler)，需指定坐标是否归一化(normalize)，地址类型(addressing mode)和滤波方式(filter mode). 在GPU中，图像对象将被映射到纹理对象上，因此会有更高的读取速度。

更多细节可参考OpenCL Specification.
代码实例可参考NVIDIA或者AMD发布的SDK中的例子。

Tags: concept, OpenCL, Tutorial |  Print | 952 views