原文地址: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将无法看见。