kernel - OpenCL-为什么使用READ_ONLY或WRITE_ONLY缓冲区

Question

在OpenCL中，将缓冲区标记为READ_ONLY或WRITE_ONLY有什么性能上的好处？

这是我经常看到的kernel(a是READ_ONLY和b是WRITE_ONLY):

__kernel void two_buffer_double(__global float* a, __global float* b)
{
    int i = get_global_id(0);
    b[i] = a[i] * 2;
}

此 kernel似乎更好，因为它使用较少的全局内存(a为 READ_WRITE):

__kernel void one_buffer_double(__global float* a)
{
    int i = get_global_id(0);
    a[i] = a[i] * 2;
}

是否仅存在 READ_ONLY和 WRITE_ONLY标志可帮助调试和捕获错误？

Answer 1

为了直接回答您的问题，我要说:不，这些标志不仅仅用于帮助调试和捕获错误。但是，很难给出任何实现如何使用这些标志以及它们如何影响性能的引用。

我的理解(很遗憾，没有任何文档支持)是，当使用这些标志时，您对缓冲区的使用方式施加了更多限制，因此可以帮助运行时/驱动程序/编译器做出一些可能改善性能的假设。 。例如，我想像那样，内核使用它时，应该不担心只读缓冲区的内存一致性，因为不应将工作项写入其中。因此可以跳过某些检查...尽管在Opencl中，您应该自己使用障碍物等来解决这一问题。

还应注意，自Ope​​ncl 1.2起，这次引入了一些其他标志，这与主机如何访问缓冲区有关。有:

CL_MEM_HOST_NO_ACCESS,
CL_MEM_HOST_{READ, WRITE}_ONLY,
CL_MEM_{USE, ALLOC, COPY}_HOST_PTR

我猜再次它必须帮助实现opencl的人员提高性能，但是我想我们需要一些AMD或NVIDIA专家的意见。

请注意，到目前为止，我所说的只是 我的想法，而不是基于任何严肃的文档(我没有找到任何文档)。

另一方面，我可以肯定地告诉您，该标准不会像@Quonux所说的那样将只读缓冲区强制置于恒定空间中。可能某些实现为小型缓冲区执行了此操作。别忘了常量空间内存很小，因此您可以读取太大而无法容纳的缓冲区。确保缓冲区位于常量空间内存中的唯一方法是在内核代码中使用常量关键字，如下所示:解释 here。当然，在主机端，如果要使用常量缓冲区，则必须使用只读标志。