Metal:没有调用内核中的“threadgroup_barrier”或使用未声明的标识符'mem_threadgroup‘的匹配函数

Question

我正在使用苹果的金属编写一个计算函数(即内核)来进行科学计算。

在内核中，我使用螺纹组内存空间。(据我理解，它类似于OpenCL中的本地内存空间--如果我错了，请纠正我。)为了同步一些内存读/写操作，我需要放置threadgroup_barrier(mem_threadgroup)。但是，屏障命令始终产生错误：

Use of undeclared identifier 'mem_threadgroup.' 

即使当我移除函数调用(threadgroup_barrier())的参数时，也会得到一个错误：

No matching function for call to 'threadgroup_barrier' in the kernel.

我在内核中包含了标题'metal_stdlib‘。我在这里错过了什么？我还需要用另一个标题来使用屏障吗？

如有任何建议，将不胜感激。

以下是代码摘要：

#include <metal_stdlib>
using namespace metal;

kernel void myKernel(device float2 *args [[buffer(0)]],
                    uint2 bidx [[threadgroup_position_in_grid]],
                    uint2 tidx [[thread_position_in_threadgroup]])
{
    // memory space shared by thread groups
    threadgroup float2 tile[32][32+1];

    ...

    for (uint k = 0; k < params.depth; k++)
    {
       ... // operations with tile (threadgroup memory space)

       threadgroup_barrier(mem_threadgroup);

       ... // more operations with tile

       threadgroup_barrier(mem_threadgroup);
    }
}

Answer 1

多亏了我的同事，他帮我找到了修复方法。因为mem_flags是一个枚举类，所以我需要范围解析操作符(mem_flags::)。因此，正确使用屏障是

threadgroup_barrier(mem_flags::mem_threadgroup)

这个修正消除了错误。