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

这取决于，

一个READ_ONLY __global内存位置存储在“全球/常量内存数据缓存”，这是很多比在GPU（看到正常的高速缓存或内存更快这里），不要紧在CPU上。

我不知道 WRITE_ONLY 的任何优点，也许它也有帮助，因为 GPU 知道它可以在不需要缓存的情况下将数据流式传输。

如果您不确定，请去测量它...

Answer 2

为了直接回答您的问题，我会说：不，这些标志的存在不仅仅是为了帮助调试和捕获错误。然而，很难就任何实现如何使用这些标志以及它们如何影响性能提供任何参考。

我的理解（不幸的是没有任何文档支持）是，在使用这些标志时，您对缓冲区的使用方式施加了更多限制，因此您可以帮助运行时/驱动程序/编译器做出一些可能提高性能的假设。例如，我认为在内核使用只读缓冲区时不应该担心内存一致性，因为工作项不应该写入其中。因此，可以跳过一些检查……尽管在 Opencl 中，您应该使用障碍等自己处理这个问题。

另请注意，自 Opencl 1.2 以来，这次引入了一些其他标志，与主机需要如何访问缓冲区相关。有：

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

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

请注意，到目前为止我所说的只是我的想法，并不是基于任何严肃的文件（我没有设法找到任何文件）。

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

Answer 3

注意，实际上有两种。您有CL_MEM_READ_ONLY，CL_MEM_WRITE_ONLY并且CL_MEM_READ_WRITE在分配缓冲区时也有__read_only，__write_only并且还具有，并__read_write用来装饰内核代码中的指针。

这些可用于优化和错误检查。让我们先来看一下性能。如果遇到只写缓冲区，则不需要缓存写入（如通过高速缓存写入），从而为读取节省更多的缓存。这在很大程度上取决于GPU硬件，至少NVIDIA硬件确实具有实际实现此功能所需的说明（.csand .lu修饰符）。您可以参考他们的PTX ISA。我没有看到编译器实际执行此优化的任何证据，例如：

__kernel void Memset4(__global __write_only unsigned int *p_dest,
    const unsigned int n_dword_num)
{
    unsigned int i = get_global_id(0);
    if(i < n_dword_num)
        p_dest[i] = 0; // this
}

被编译为：

st.global.u32 [%r10], %r11; // no cache operation specified

这是有道理的，因为CUDA没有这些限定符的等效项，因此编译器很可能会静默地忽略那些限定符。但是将它们放置在这里并没有什么坏处，将来我们可能会更幸运。在CUDA中，某些功能是通过使用该__ldg功能以及使用编译器标志来选择启用/退出对L1（-Xptxas -dlcm=cg）中的全局内存传输进行缓存的公开。asm如果您发现绕过缓存可以带来很大的好处，也可以随时使用。

至于错误检查，使用const内核声明中的指定符可以很容易地避免写入只读缓冲区。在纯“ C”中，不允许从只写缓冲区中读取数据是不可能的。

将这些缓冲区映射到主机内存时，可能会发生另一种优化。映射CL_MEM_READ_ONLY缓冲区时，映射的区域可能会保留为未初始化，因为主机只会写入该内存，而设备只能读取该内存。同样，在取消映射CL_MEM_WRITE_ONLY缓冲区时，驱动程序不需要将（可能由主机修改的）内容从主机内存复制到设备内存。我没有测量。

作为附带说明，我尝试使用：

st.global.u32 [%r10], %r11; // no cache operation specified

即使在带有sm_35设备的简单memcpy内核上（在GTX 780和K40上测试），它也可以为您提供约15 GB /秒的额外存储空间。尚未看到明显的加速效果sm_30（不确定是否从ptx中删除了指令，但是否在那里甚至还希望它得到支持）。请注意，您需要定义NVIDIA自己（或参见内核代码中的“检测OpenCL设备供应商”）。