OpenCL全局同步尝试失效：内核执行异常及死锁问题排查

让我们一步步拆解你代码里的问题，以及为什么会出现这些奇怪的行为：

1. 循环仅执行一次就退出的核心原因 #

你的g_barrier变量在第一次循环后没有被重置！

第一次循环时，i=1，t = 1 * get_num_groups(0)。每个工作组的第0个线程执行atomic_add(g_barrier, 1)，当所有工作组都完成这一步后，g_barrier的值会等于get_num_groups(0)，满足*g_barrier >= t，退出while循环，进入下一次i=2的循环。

但此时g_barrier的值还是停留在get_num_groups(0)，而新的t = 2 * get_num_groups(0)，这时候*g_barrier < t的条件一开始就不成立，while循环直接跳过，后续所有循环都会重复这个逻辑，导致看起来内核只执行了一次循环。

2. 添加volatile后死锁的原因 #

当你把g_barrier声明为volatile __global uint *时，虽然强制每次读取都从全局内存取（避免缓存优化导致的读取旧值），但你依然没解决屏障未重置的核心问题：

• 第一次循环结束后，g_barrier停留在get_num_groups(0)；
• 第二次循环时t=2*get_num_groups(0)，所有线程都会卡在while(*g_barrier < t)里——因为没有任何代码会把g_barrier重置为0并重新累加，最终导致内核无限挂起，系统资源被占满，出现黑屏卡顿和死锁。

另外，自定义全局屏障本身就有风险：不同设备的全局内存缓存一致性模型不同，即使你用了atomic_add和volatile，也无法保证所有工作组能同时看到g_barrier的最新值，这会进一步加剧同步问题。

正确的全局同步实现方式 #

如果你需要跨所有工作组的全局同步，推荐以下两种方案：

方案1：在主机端同步（OpenCL 1.x兼容） #

OpenCL 1.x没有原生的内核内全局同步机制，最可靠的方式是把循环拆到主机端，每次执行完一轮内核后，在主机端调用同步接口，再启动下一轮内核：

// 主机端代码示例
for (int i = 1; i < MAX; i++) {
    // 设置本轮内核需要的参数（比如当前i的值）
    clSetKernelArg(kernel, ...);
    // 启动内核执行本轮计算
    clEnqueueNDRangeKernel(command_queue, kernel, ...);
    // 等待所有工作组完成本轮计算，确保全局状态同步
    clFinish(command_queue);
}

这种方式完全依赖OpenCL的主机端同步，避免了内核内自定义全局屏障的各种坑。

方案2：使用OpenCL 2.0的原生全局同步（设备支持时） #

如果你的设备支持OpenCL 2.0，可以结合work_group_barrier和原子操作实现内核内全局同步，但必须做好屏障重置：

__kernel void Some_Kernel(__global uint *g_barrier) {
    uint i, t;
    for (i = 1; i < MAX; i++) {
        // 你的业务计算代码
        // 先确保工作组内所有线程完成当前计算
        work_group_barrier(CLK_GLOBAL_MEM_FENCE);
        
        // 仅每个工作组的第0个线程参与全局屏障计数
        if (get_local_id(0) == 0) {
            atomic_add(g_barrier, 1);
        }
        // 等待所有工作组完成计数
        work_group_barrier(CLK_GLOBAL_MEM_FENCE);
        
        t = i * get_num_groups(0);
        while(atomic_load_explicit(g_barrier, memory_order_relaxed) < t) {
            // 空循环等待，用relaxed原子读取优化性能
        }
        
        // 关键：所有线程完成同步后，重置屏障（仅一个线程执行即可）
        work_group_barrier(CLK_GLOBAL_MEM_FENCE);
        if (get_group_id(0) == 0 && get_local_id(0) == 0) {
            *g_barrier = 0;
        }
        work_group_barrier(CLK_GLOBAL_MEM_FENCE);
    }
}

这里必须在每次循环结束后重置g_barrier为0，同时用work_group_barrier确保所有线程能看到重置后的最新值。

内容的提问来源于stack exchange，提问作者sdml