Metal：内核中没有与“threadgroup_barrier”匹配的函数或使用未声明的标识符“mem_threadgroup”。

Question

我正在使用苹果的Metal进行科学计算，编写一个计算函数（也称为内核）。

在这个内核中，我使用了threadgroup内存空间。（据我理解，它类似于OpenCL中的local内存空间 - 如果我错了，请纠正我。） 为了同步一些内存读/写操作，我需要放置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);
    }
}

- JustDoIt

Answer 1

[感谢我的同事帮助找到修复方法。] 由于mem_flags是一个枚举类，我需要使用作用域解析运算符（mem_flags::）。因此，屏障的正确用法是

threadgroup_barrier(mem_flags::mem_threadgroup)

这个修复程序消除了错误。

- JustDoIt