OpenCL 工作组

如之前类比学校的例子，工作项除了在年级中有ID（全局ID），在班级（工作组）中也有ID。工作组在工作项需要同步时显得十分重要，同时对于局部存储器是以工作组为个体来分配的，工作组内的工作项可以共享局部存储器。在需要使用局部存储器来优化存储器访问的应用中，我们需要知道工作项在工作组内的ID，从而使我们可以更好地控制每个工作项执行的数据。在OpenCL中提供了这些函数来查询工作组的信息，以及工作项在工作组内的ID，详细见下表。

表格中的几个函数，意思都很直观。下面我们通过一个比较复杂的例子来包含本节中讲解的几个我们常用的工作项布局函数，代码如下：

主机端：

……
const size_t globalSize[2] = {6, 4};
const size_t localSize[2] = {3, 2}
const size_t offset[2] = {3, 5};
size_t dim = 2;
err = clEnqueueNDRangeKernel(cmdqueue, kernel, dim, offset,
                                  globalSize, localSize, 0, NULL, NULL);
……

内核代码：

kernel void id_check(global float ＊output)
{
    size_t global_id_0 = get_global_id(0);
    size_t global_id_1 = get_global_id(1);
    size_t global_size_0 = get_global_size(0);
size_t offset_0 = get_global_offset(0);
size_t offset_1 = get_global_offset(1);
size_t local_id_0 = get_local_id(0);
size_t local_id_1 = get_local_id(1);
int index_0 = global_id_0 - offset_0;
int index_1 = global_id_1 - offfset_1;
int index = index_1 ＊ global_size_0 + index_0;
float f = global_id_0 ＊ 10.0f + global_id_1 ＊ 1.0f;
f += local_id_0 ＊ 0.1f + local_id_1 ＊ 0.01f;
output[index] = f;
}

如上代码，output的输出结果为：

35.00      45.10          55.20          65.00          75.10          85.20
36.01      46.11          56.21          66.01          76.11          86.21
37.00      47.10          57.20          67.00          77.10          87.20
38.01      48.11          58.21          68.01          78.11          88.21

这个例子中，把函数clEnqueueNDRangeKernel（）中关于工作项设定的所有参数都用上，对于刚接触OpenCL的读者可能理解上有点困难。对结果不能很好理解的读者，建议在电脑上试验上述代码，直到对工作项布局函数都理解透彻为止。