话说OpenCL(一)

 

今天扫了一下OpenCL，大概印象是这样的（这里不在语法细节上纠结，只在大面儿上谈一下）：OpenCL（Open Compute Language），异构系统（CPU+GPU）并行编程的一个开放标准，接口比较底层，可视为NV自家API向开放标准的一次升级或叫过渡。主要元素有：kernel程序可以在运行时编译（不像CUDA driver API那样由NVCC.exe预编译成ptx/cubin再作为module进行加载），这样可以针对特定上下文特定设备在运行时进行编译；引入buffer object概念进行囊括显存、kernel加载等操作，相当于引入了一个管家婆。

具体细节，请参看下面演示程序，演示程序来自OpenCL快速入门指南：

 

【演示程序】矢量加（kernel程序）

 

图中三处标注，第三处get_global_id(0)计算得到线程在work group(OpenCL将block称为work group)使用户屏蔽掉线程号计算的麻烦，此外OpenCL提供了get_local_id() get_work_dim() get_global_size()分别用来获得block内线程号、block内线程数、block个数。

 

【演示程序】 主机端程序（加载矢量加）

a)       CUDA驱动程序，首先NVCC.exe将其预编译为CUBIN文件vectorAdd.cubin

b)      OpenCL程序，加载的是sProgramSource源程序。

 

对比这两段程序：

1）  创建上下文的时候，CUDA先选设备，再为选择的设备创建上下文；而OpenGL是为某种类型的设备（比如这里是CL_DEVICE_TYPE_GPU）创建上下文，然后寻找支持这种上下文的设备，之后进行选择；

2）  为这个设备的这个上下文创建命令队列，命令队列里放的是显存读写操作和kernel加载指令；

3）  接着要加载并编译kernel程序啦。注意，这里的编译是特定于这个上下文环境的编译。

4）  下面创建kernel。其实就是在编译后的hProgram里找到VectorAdd那段儿程序，返回一句柄。

5）  接下来分配内存空间、初始化这段空间、分配显存空间。注意，OpenCL里显存是通过buffer对象来管理的。Buffer比CUDA内存管理函数意义要更丰富，可以设置只读、只写模式，也可以设置这段内存区域可直接被设备访问。

6）  执行参数设置，执行kernel。注意，OpenCL里通过clEnqueNDRangeKernel()来加载，可以指定work group和global group的size，都可以是1D/2D/3D。

7）  数据回拷。注意，OpenCL里通过clEnqueReadBuffer()来实现，可以指定回拷方式，例如同步读或异步读。

8）  善后。

 

【几点不同】CUDA driver API  VS. OpenCL

1.       指针传递

struct Node{Node* next;}

n = n -> next;

这个在CUDA里是允许的，但在OpenCL里不可以。在OpenCL里的指针，基于buffer和offset这样一组概念，即在同一个buffer里（不同buffer不能引用）用offset来引用，比如：

struct Node {unsigned int next;}

…

n = bufBase + n;

 

2.       kernel程序的执行方式

CUDA里kernel程序是先被nvcc.exe预编译成ptx / cubin文件，然后在一个CUDA上下文里将其作为module加载并调用其中的函数。而OpenCL程序是将kernel源程序包进来，在运行时根据你选择的设备进行编译（这是不是就是一定程度的跨平台了呢？），当然也可以像CUDA那样编译成二进制再进行加载（总不至于抛弃前辈是吧）。

 

3.       kernel加载内存偏移