金属核函数是原子的吗？

Question

内核函数内部的行之间可以发生上下文切换吗？

因为我在进行更改之前设置了一些值，所以我想确定是否设置了该值，是否已经进行了更改。

Answer 1

简短的回答是，是的，上下文切换肯定会“在两行之间”发生。这就是上下文切换的重点：如果着色器中的某些行（无论是片段、顶点还是内核）需要一些尚不可用的资源（ALU、特殊功能单元、纹理单元、内存），GPU 肯定会切换上下文。这称为延迟隐藏，它对 GPU 的性能非常重要，因为没有它，GPU 内核将大部分时间花在上述不同资源上。所有这一切都意味着 Metal 核函数绝对不是原子的。

至于您遇到的问题，如果您希望某事原子地发生，则在 Metal Shading Language 中有两种主要方法：

1. 您可以使用metal_atomicheader 中的原子类型和函数。这是 C++14atomic头文件的子集，它包含原子存储、加载、交换、比较和交换以及获取和修改功能。
2. 您可以使用 SIMD 组和线程组屏障。屏障允许在允许任何线程继续之前等待组中的所有线程执行所有操作。根据传递给屏障函数的标志，屏障还可以对内存访问进行排序或仅用作执行屏障（如果您从着色器中写入一些数据，这可能不是您想要的）。

您可以参考Metal Shading Language Specification获取更多关于这些函数的信息（查看 5.8.1 节获取有关“线程组和 SIMD 组同步函数”的信息，查看 5.13 节获取有关“原子函数”的信息）。

通常，如果您的内核函数处理了一些您以后需要减少的数据，您将执行以下操作（这是一个非常简单的示例）：

kernel void
my_kernel(texture2d<half> src [[ texture(0) ]],
    texture2d<half, access::write> dst [[ texture(1) ]],
    threadgroup float *intermediate [[ threadgroup(0) ]],
    ushort2 lid [[ thread_position_in_threadgroup ]],
    ushort ti [[ thread_index_in_threadgroup ]],
    ushort2 gid [[ thread_position_in_grid ]])
{
    // Read data
    half4 clr = src.read(gid);

    // Do some work for each thread
    intermediate[ti] = intermediateResult;

    // Make sure threadhroup memory writes are visible to other threads
    threadgroup_barrier(mem_flags::mem_threadgroup);

    // One thread in the whole threadhgroup calculates some final result
    if (lid.x == 0 && lid.y == 0)
    {
        // Do some work per threadgroup
        dst.write(finalResult, gid);
    }
}

这里线程组中的所有线程从src纹理中读取数据，执行工作，将中间结果存储在线程组内存中，然后计算并将最终结果写入纹理dst。threadgroup_barrier可以确保其它线程（包括线程thread_position_in_threadgroup等于(0, 0)说是要计算最终结果），可以看到内存的写入。