OpenCL中亦要注意volatile的使用

原文地址：http://blog.csdn.net/zenny_chen/article/details/6130109

在OpenCL或CUDA中，对全局共享变量的访问我们往往会忽略掉使用volatile，这在仅对其访问一次的时候不会有问题，但是对这个共享变量变量做第二次访问的话，那么它会被编译器优化，从而得到的一直是第一次被引用时的值。也就是说，其它线程对共享变量的修改，当前线程将不可见。

下面举一个OpenCL的简单的例子来描述这一情况：

__kernel void solve_sum(
                    __global unsigned buffer[512],
                    __global unsigned dest[512]
                    )
{
    __local volatile int flag = 0;
    
    size_t gid = get_global_id(0);
    
    const uint4 value = (uint4)(1, 2, 3, 4);
    
    if(0 <= gid && gid < 32)
    {
        while(flag == 0);
        vstore4(value, gid, buffer);
        //write_mem_fence(CLK_GLOBAL_MEM_FENCE);
        flag = 0;
    }
    else if(32 <= gid && gid < 64)
    {
        flag = 1;
        while(flag == 1);
        unsigned ret = buffer[127 + 32 - gid];
        
        dest[gid - 32] = ret;
    }
}

在以上代码中，如果把volatile去掉，那么线程32到63这一warp将处于死循环。由于之前对flag写了1，因此在后面while(flag == 1);这句执行时，将一直为true；外部对flag的修改，此warp将无法看见。