8

在中，将缓冲区标记为或OpenCL是否有任何性能优势？READ_ONLYWRITE_ONLY

这kernel是我经常看到的（a is READ_ONLYand b is 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 is 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标志只是为了帮助调试和捕获错误而存在吗？

4

3 回答 3

7

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

这些可用于优化和错误检查。让我们先看看性能。如果遇到只写缓冲区，则无需缓存写入（如通过缓存写入），为读取节省更多缓存。这在很大程度上取决于 GPU 硬件，至少 NVIDIA 硬件确实具有实际实现它所需的指令（.cs和.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缓冲区时，驱动程序不需要将（可能被主机修改的）内容从主机内存复制到设备内存。我没有测量这个。

作为旁注，我尝试使用：

inline unsigned int n_StreamingLoad(__global __read_only const unsigned int *p_src)
{
#ifdef NVIDIA
    unsigned int n_result;
    asm("ld.global.cs.u32 %r0, [%r1];" : "=r" (n_result) : "r" (p_src));
    return n_result;
#else // NVIDIA
    return *p_src; // generic
#endif // NVIDIA
}

inline void StreamingWrite(__global __write_only unsigned int *p_dest, const unsigned int n_value)
{
#ifdef NVIDIA
    asm("st.global.cs.u32 [%r0], %r1;" : : "r" (p_dest), "r" (n_value) : "memory");
#else // NVIDIA
    *p_dest = n_value; // generic
#endif // NVIDIA
}

即使在带有sm_35设备的简单 memcpy 内核上（在 GTX 780 和 K40 上测试），它也可以为您提供大约 15 GB/秒的额外速度。还没有看到明显的加速sm_30（不确定它是否甚至意味着在那里得到支持——尽管指令没有从 ptx 中删除）。请注意，您需要NVIDIA自己定义（或参见在内核代码中检测 OpenCL 设备供应商）。

于 2015-12-10T14:30:49.367 回答
5

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

我的理解（不幸的是没有任何文档支持）是，当使用这些标志时，您会对缓冲区的使用方式施加更多限制，因此您可以帮助运行时/驱动程序/编译器做出一些可能提高性能的假设。例如，我想在内核使用只读缓冲区时，不应该担心内存一致性问题，因为工作项不应该写入其中。因此，可以跳过一些检查……尽管在 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 所说的那样强制只读缓冲区位于常量空间中。可能是某些实现对小缓冲区执行此操作。我们不要忘记常量空间内存很小，因此您可以将只读缓冲区太大而无法放入。确保缓冲区位于常量空间内存中的唯一方法是在内核代码中使用常量关键字作为在这里解释。当然在主机端，如果你想使用常量缓冲区，你必须使用只读标志。

于 2013-07-27T21:21:36.897 回答
5

这取决于，

READ_ONLY __global内存位置存储在“全局/常量内存数据缓存”中，这比GPU 上的普通缓存或 RAM 快得多（请参阅此处），在 CPU 上没关系。

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

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

于 2013-07-27T19:31:07.843 回答