Apple's OpenCL——多线程同步
http://www.cocoachina.com/bbs/read.php?tid=37608
我们前几章介绍了OpenCL的一些基本概念以及一些基本的用法。我们之前的例子都是线程独立计算的，相互之间没有任何通信。而这样的计算模型也是GPU最最喜欢的，能完全发挥GPU众核并行计算的优势。 今天我们将介绍OpenCL多线程同步技巧。我们下面的例子将是以一个简单的求和算法来描述如何同步一个工作组内的线程以及工作组之间如何同步。 我们之前介绍过变量的地址属性。用__global修饰的变量存放在显示存储器中，特点是容量很大，但访问速度很慢，并且所有工作项都能访问；而用__local修饰的变量存放在共享存储器，其特点是速度比全局存储要快很多，并且在同一工作组内的工作项能够对其进行访问，而且每个工作组有自己独立的共享存储器；__private修饰或默认状态下定义的变量是私有的，即存放在寄存器中，其特点是访问速度相当快，基本上一次读或写仅需要1个着色器周期，但它是工作项私有的，并且每个工作项只有若干个寄存器可以进行访问。 如果我们让在一个工作组内的线程进行同步，那么我们可以借助共享存储变量来帮我们达成这个目标；而如果是工作组之间的通信，则需要全局存储变量。 下面看求和的内核代码：
__kernel void solve_sum(
__global int input[4096],
__global int output[9]
)
{
__local int localBuffer[512];
size_t item_id = get_local_id(0);
size_t gid = get_global_id(0);
localBuffer[item_id] = input[gid];
barrier(CLK_LOCAL_MEM_FENCE);
if((item_id) == 0)
{
int s = 0;
for(int i = 0; i < 512; i++)
s += localBuffer[ i ];
output[get_group_id(0)] = s;
output[8] = get_num_groups(0);
}
}
在以上代码中，一共有4096个工作项，共有8个工作组，这样每个工作组就有512个工作项。这个算法很简单，首先将每个工作组内的工作项存放到共享数组中，等到一个工作组内的所有工作项完成这个动作后，让工作项0对共享存储缓存中的数据进行求和，完成后写入到相应的工作组索引的输出缓存。
在上述代码中，get_local_id获得的是当前工作组中的当前工作项索引，在上述代码环境中的范围是0到511。 因此，我们可以将localBuffer[item_id] = input[gid];这句改为：localBuffer[gid & 511] = input[gid];这两条语句的语义完全等价。 这里要着重介绍的线程同步函数是：
void barrier (cl_mem_fence_flags flags)
这个内建函数对应于处理器的一条指令，其作用是同步一个工作组内的所有工作项。我们现在把工作项看作为一个线程。当其中一个线程执行到barrier时，它会被处理器阻塞住，直到该工作组内所有线程都执行到这个barrier，然后这些线程才能继续执行下去。 这里有一个参数flags用于指示存储器栅栏是局部的还是全局的，我们这里只需要局部的，因为这里不需要工作组之间的同步。 我们把每个工作组计算出来的结果写到输出缓存中。由于输出才8个32位数据，因此在CPU中再拿去计算也变成了小菜一碟，呵呵。 下面附上整个工程的代码。
附件: 
OpenCL_Basic.zip (17 K) 下载次数:116
上述代码是将每个工作组计算好的结果传送给主机端。那么我们是否能让GPU把这8个结果也一起解决掉呢？答案是肯定的。 不过我们这里将会用到OpenCL1.0中的原子操作扩展。这些基于int32位的原子操作在OpenCL1.1中将正式归为语言核心，而不是扩展。 我们可以通过OpenCL查询获得 cl_khr_global_int32_base_atomics是否被支持。如果被支持，那么我们可以用下面的方法：
__kernel void solve_sum(
__global int input[4096],
__global int output[9]
)
{
__local int localBuffer[512];
size_t item_id = get_local_id(0);
size_t gid = get_global_id(0);
localBuffer[item_id] = input[gid];
barrier(CLK_LOCAL_MEM_FENCE);
if(item_id == 0)
{
int s = 0;
for(int i = 0; i < 512; i++)
s += localBuffer[i];
output[get_group_id(0)] = s;
int index = atom_inc(&output[8]);
if(index == 7)
{
mem_fence(CLK_GLOBAL_MEM_FENCE);
s = 0;
for(index = 0; index < 8; index++)
s += output[index];
output[8] = s;
}
}
}
在上述代码中，我们用了原子累积操作：
int atom_inc (__global int *p)
这个函数是先读取p指针所指地址的内容，然后将该内容递增1，最后写回到这个地址中去，并且返回读到的那个值（即更新以前的值）。整个操作都是不被打断的，因此是一个原子操作。 我们在上述代码中，用一个索引来获取返回值，如果索引为7，说明当前线程是最后一个写结果的工作组中的第0个线程。于是，我们利用这个线程把8个结果累加，然后写回到输出缓存。 如果有两个线程对同一地址同时执行atom_inc，那么GPU将会进行仲裁，它只允许其中一个执行这一操作，而等到这个操作完成之后，其它线程才能继续，否则，其它要执行此操作的线程都将被处理器阻塞。
那么这里由于利用了输出缓存作为全局存储的计数器变量，因此它将不象第一份代码那样作为只写参数，而是要设置为可读可写的参数，并且要把初始数据传入给GPU设备端。下面附上相应的工程和代码。
附件: OpenCL_Basic.zip (17 K) 下载次数:53
第27行显式地使用了memory fence主要是为了保证其它所有的工作组把各自计算得到的结果都完全写回到全局的输出缓存中。 关于memory fence的概念以及作用可以参考：Memory Fence 当然，这里把它省略掉也完全没有问题，呵呵。
下面要讲一下关于Local Memory的一些高级话题。
什么是共享存储器段呢？一个共享存储器段就是在共享存储器中的一个32位字（当前主流的中低端GPU均是如此，高级点的则可能是64位或更大）。那么，如果一个工作组的共享存储器空间是128KB的话，则共有128KB / 4B = 32 * 1024个段。 如果有两个线程（即工作项）对同一个段进行写操作，那么这些写操作将由原来可以并行写而变成串行化的写，也就是说，总线控制器会对这些多个线程的写进行串行化，它会选择其中一个线程先写，完了之后再挑选下一个。那么这样一来，多个线程的执行也就从原来的并行操作变成了串行操作，这样会受到很大的性能惩罚。 因此，我们在设计算法时应该尽量保证每个线程只对自己相应的共享存储器段进行写操作，而避免有多个线程去写同一个共享存储器段。 而像上面示例代码中，由于读写的数据元素都是32位，正好是一个存储器段的大小，并且一个工作组内的每个工作项都以自己id作为索引对共享存储器进行写，这样每个工作项所写的段都是相互独立的，因此这里不会发生段冲突。
呵呵，不知道各位对Map-Reduce了解多少。上述对n项元素求和，在OpenCL的实现中就与Map-Reduce非常相似。首先将所有元素映射到共享缓存中，这个过程是Map；然后利用一个线程将本工作组内的所有元素进行求和，这个过程就是Reduce，即，将多项元素缩减到更少的元素。在图形架构中有个专用名词叫Reduction，其实就是指将向量缩减为标量。
