OpenCL - 为什么使用只读或只写缓冲区

Question

OpenCL - 为什么使用只读或只写缓冲区

9

在OpenCL中，将缓冲区标记为“READ_ONLY”或“WRITE_ONLY”是否具有性能优势？

这是我经常看到的内核（其中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标志只是为了帮助调试和捕获错误吗？

- benshope

3个回答

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所说的那样。对于小缓冲区，某些实现可能会这样做。不要忘记常数空间内存很小，因此您可能会有太大而无法放入常数空间中的只读缓冲区。确保缓冲区位于常数空间内存中的唯一方法是在内核代码中使用constant关键字，如此处所述。当然，在主机端，如果您想使用常量缓冲区，则必须使用只读标志。

- CaptainObvious

5

这要看情况而定。

READ_ONLY __global 内存位置存储在“全局/常量内存数据缓存”中，这比GPU上的普通缓存或RAM快得多（请参见这里），但对于CPU而言则无关紧要。

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

如果你不确定，就去测量一下吧...

- Quonux

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- the swine · Accepted Answer

请注意，实际上有两种类型。在分配缓冲区时，您有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内核操作时，也可以提供约15个额外的GB / sec（已在GTX 780和K40上进行了测试）。在sm_30上没有看到明显的加速（不确定甚至是否支持该版本-虽然指令未从ptx中删除）。请注意，您需要自己定义NVIDIA（或参见在内核代码中检测OpenCL设备供应商）。