OpenCL 的执行模型

OpenCL 程序的运行模式
使用 OpenCL进行并行计算的首要任务是分析算法的并行性，然后将算法映射到 OpenCL的执行模型（运行模式）上。OpenCL程序的运行模式可以概括为以下三点：
1）在 OpenCL中，需要由用户使用 OpenCL C语言编写实现特定算法的 kernel（核函数），并利用运行时环境将 kernel分发到某个计算单元（例如 GPU）上运行；
2）kernel启动后会创建大量的线程同时执行，每个线程称为工作单元（work-item），用来执行kernel程序；
3）当映射到 OpenCL硬件上执行时，采用两级并行机制，工作组（work-group）并发运行在异构计算设备的计算单元（compute unit）上，同一个工作组（work-group）里的多个工作单元（work-item）在处理单元（process element）上相互独立并行执行。

在 OpenCL， 多个 work-item可以组织成一个 work-group，多个 work-group可以组织成一个 NDRange。work-group和 NDRange都可以是 1，2，3维，分别对应 x，y，z维，例如一个 work-group的大小可以是 x维有 4个线程，y维有 3个线程，z维有 2个线程，一个 NDRange的大小可以是 x维有 8个线程，y维有 6个线程，z维有 4个线程，那么这个 NDRange可以由 8个上述大小的 work-group组成。

SIMT 计算模型
这些 work-item之所以能够并行执行，是因为在硬件上采用了 SIMT（Single Instruction Multiple Threads，单指令多线程）计算模型，在 SIMT中，多个线程同时执行同一条指令，因此它们具有相同的 PC值（程序计数器），当出现分支指令时，通过 谓词寄存器 控制一部分线程执行后续指令，而另一部分线程不执行后续指令。

数据与 work-item 的映射
在 SIMT模型中，多个 work-item执行的是同一条指令，如果他们处理的数据也是同一条的话，那多线程就没有意义了。因此在核函数的编写中，程序员需要显式的将数据分配给不同的 work-item进行处理。这其中用到的关键技术就是各个 work-item相互独立的索引变量。

在 OpenCL中，
1）可以通过 get_global_id(dim)获得当前 work-item在它所属的 NDRange中 dim维的索引；
2）可以通过 get_local_id(dim)获得当前 work-item在它所属的 work-group中 dim维的索引；
3）可以通过 get_global_size(dim)获得当前NDRange在 dim维的 work-item数量；
4）可以通过 get_local_size(dim)获得当前 work-group在 dim维的 work-item数量；
不一而足，具体可以参考 OpenCL C语言规范获得更多的设备侧 API信息。

下面进行举例说明，将一个 uint类型的数组的每个元素自增 1，如果数组元素比较少，而使用的 work-item比较多，则可以每个线程只处理一个数据，如下所示：
1）线程和数据一对一

__kernel void inc(__global uint* ptr, int n) {
	int idx = get_global_id(0);
	if (idx < n)
	  ptr[idx] += 1;
}

如果数组元素比较多，而使用的 work-item比较少，则一个线程需要处理多个数据，如下所示：
2）线程和数据一对多

__kernel void inc(__global uint* ptr, int n) {
	int size = get_global_size(0);
	int idx = get_global_id(0);
	for (; idx < n; idx += size) {
	  ptr[idx] += 1;
	}
}