OpenCL 的架构
OpenCL 包括一组 API 和一个程式语言。基本上,程式透过 OpenCL API 取得 OpenCL 装置(例如显示晶片)的相关资料,并将要在装置上执行的程式(使用 OpenCL 程式语言撰写)编译成适合的格式,在装置上执行。OpenCL API 也提供许多装置控制方面的动作,例如在 OpenCL 装置上取得一块记忆体、把资料从主记忆体复制到 OpenCL 装置上(或从 OpenCL 装置上复制到主记忆体中)、取得装置动作的资讯(例如上一个程式执行所花费的时间)等等。
例如,我们先考虑一个简单的工作:把一群数字相加。在一般的 C 程式中,可能是如下:
float a[DATA_SIZE];
float b[DATA_SIZE];
float result[DATA_SIZE];
// ...
for(int i = 0; i < DATA_SIZE; i++) {
result[i] = a[i] + b[i];
}
在 OpenCL 中,则大致的流程是:
- 把 OpenCL 装置初始化。
- 在 OpenCL 装置上配置三块记忆体,以存放 a、b、c 三个阵列的资料。
- 把 a 阵列和 b 阵列的內容,复制到 OpenCL 装置上。
- 编译要执行的 OpenCL 程式(称为 kernel)。
- 执行编译好的 kernel。
- 把计算結果从 OpenCL 装置上,复制到 result 阵列中。
透过 data parallel 的模式,这里的 OpenCL 程式非常简单,如下所示:
__kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
int idx = get_global_id(0);
result[idx] = a[idx] + b[idx];
}
在一般的版本中,是透过一个回圈,执行 DATA_SIZE
次数的加法动作。而在 OpenCL 中,则是建立DATA_SIZE 个
work item,每个 work item 都执行上面所示的 kernel。可以看到,OpenCL 程式语言和一般的 C 语言非常类似。__kernel
表示这个函式是在 OpenCL 装置上执行的。
__global
则表示这个指标是在 global memory 中(即 OpenCL 装置上的主要記记忆体)。而
get_global_id(0)
会传回 work item 的编号,例如,如果有 1024 个 work item,则编号会分別是 0 ~ 1023(实际上编号可以是二维或三维,但在这里先只考虑一维的情形)。