运维开发网

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

运维开发网 https://www.qedev.com 2020-05-23 12:25 出处:网络 作者:运维开发网整理
在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
在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;
}

这个内核似乎更好,因为它使用较少的全局内存(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标志以帮助调试和捕获错误?

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

我的理解(遗憾的是没有任何文档备份)是,当使用这些标志时,你会对缓冲区的使用方式施加更多限制,因此你可以帮助运行时/驱动程序/编译器做出一些可能改善性能的假设.例如,我想在内核使用它时不应该担心只读缓冲区的内存一致性,因为工作项不应该写入其中.因此可以跳过一些检查……虽然在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.当然在主机端,如果你想使用常量缓冲区,你必须使用只读标志.

0

精彩评论

暂无评论...
验证码 换一张
取 消