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将无法看见。