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;
}
}