OpenCL中亦要注意volatile的使用

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