"threadgroup_barrier" 没有区别

Question

目前我正在使用 Metal 计算着色器并试图了解 GPU 线程同步在那里的工作原理。

我写了一个简单的代码，但它没有按我预期的方式工作：

假设我有 threadgroup 变量，它是一个数组，所有线程都可以同时产生一个输出。

    kernel void compute_features(device float output [[ buffer(0) ]],
                                 ushort2 group_pos [[ threadgroup_position_in_grid ]],
                                 ushort2 thread_pos [[ thread_position_in_threadgroup]],
                                 ushort tid [[ thread_index_in_threadgroup ]])
    {     
        threadgroup short blockIndices[288];

        float someValue = 0.0
        // doing some work here which fills someValue...

        blockIndices[thread_pos.y * THREAD_COUNT_X + thread_pos.x] = someValue;

        //wait when all threads are done with calculations
        threadgroup_barrier(mem_flags::mem_none);  
        output += blockIndices[thread_pos.y * THREAD_COUNT_X + thread_pos.x]; // filling out output variable with threads calculations
    }

上面的代码不起作用。输出变量不包含所有线程计算，它仅包含来自线程的值，该值可能是最后一个将值加到 output。对我来说，threadgroup_barrier 似乎什么都没做。

现在，有趣的部分。下面的代码有效：

blockIndices[thread_pos.y * THREAD_COUNT_X + thread_pos.x] = someValue;

threadgroup_barrier(mem_flags::mem_none);  //wait when all threads are done with calculations
if (tid == 0) {
    for (int i = 0; i < 288; i ++) {
        output += blockIndices[i]; // filling out output variable with threads calculations
    }
}

而且此代码也与前一个代码一样有效：

blockIndices[thread_pos.y * THREAD_COUNT_X + thread_pos.x] = someValue;

if (tid == 0) {
    for (int i = 0; i < 288; i ++) {
        output += blockIndices[i]; // filling out output variable with threads calculations
    }
}

总而言之：只有当我在一个 GPU 线程中处理线程组内存时，我的代码才能按预期工作，无论它的 ID 是什么，它都可以是线程组中的最后一个线程，也可以是第一个线程。 threadgroup_barrier 的存在绝对没有区别。我还使用了 threadgroup_barrier 和 mem_threadgroup 标志，代码仍然不起作用。

我知道我可能遗漏了一些非常重要的细节，如果有人能指出我的错误，我会很高兴。提前致谢！

Answer 1

当你写output += blockIndices[...]时，所有线程都会同时尝试执行这个操作。但是由于 output 不是原子变量，这会导致竞争条件。这不是线程安全操作。

您的第二个解决方案是正确的。您只需要一个线程来收集结果（尽管您也可以将其拆分为多个线程）。如果你移除障碍它仍然可以正常工作可能只是因为运气。

"threadgroup_barrier" 没有区别

"threadgroup_barrier" makes no difference

multithreading

gpu

threadgroup

compute-shader

metal