苹果开发语言OpenCL 多线程同步附源码-ios开发多线程

苹果开发语言OpenCL 多线程同步附源码是本文要介绍的内容，首先我们先来了解一下OpenCL，即：Open Computing Language，是由苹果公司起草设计的用于大规模并行计算的计算编程语言。

今天我们将介绍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;  
        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)

上述代码是将每个工作组计算好的结果传送给主机端。那么我们是否能让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)

下面要讲一下关于Local Memory的一些高级话题。

其实OpenCL中的local memory对应于CUDA中的shared memory。在访问共享存储器时，如果多个线程写同一个共享存储器段（memory bank），那么会导致段冲突（bank conflict）。

什么是共享存储器段呢？一个共享存储器段就是在共享存储器中的一个32位字（当前主流的中低端GPU均是如此，高级点的则可能是64位或更大）。那么，如果一个工作组的共享存储器空间是128KB的话，则共有128KB / 4B = 32 * 1024个段。

如果有两个线程（即工作项）对同一个段进行写操作，那么这些写操作将由原来可以并行写而变成串行化的写，也就是说，总线控制器会对这些多个线程的写进行串行化，它会选择其中一个线程先写，完了之后再挑选下一个。那么这样一来，多个线程的执行也就从原来的并行操作变成了串行操作，这样会受到很大的性能惩罚。

因此，我们在设计算法时应该尽量保证每个线程只对自己相应的共享存储器段进行写操作，而避免有多个线程去写同一个共享存储器段。而像上面示例代码中，由于读写的数据元素都是32位，正好是一个存储器段的大小，并且一个工作组内的每个工作项都以自己id作为索引对共享存储器进行写，这样每个工作项所写的段都是相互独立的，因此这里不会发生段冲突。

小结：苹果开发语言 OpenCL 多线程同步附源码的内容介绍完了，希望本文对你有所帮助！

帖子地址 http://www.cocoachina.com/bbs/read.php?tid-37608.html，欢迎参与讨论

苹果开发语言OpenCL 多线程同步 附源码

苹果开发语言OpenCL 多线程同步附源码