OpenCL - 缓冲区和全局内存之间的区别

Question

在 Opencl 中，缓冲区是数据与主机应用程序通信的管道。

cl_mem clCreateBuffer (cl_context context, cl_mem_flags flags, size_t size,
                       void *host_ptr, cl_int *errcode_ret);

现在如果我有一个缓冲区a_buffer 标记为READ_ONLY，并且内核是：

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

我的问题是：a_buffer 是全局内存还是常量内存？我应该为a 使用__constant 限定符吗？ cl_mem_flags(READ_ONLY和READ_WRITE)和内存限定符(global和constant)有什么联系？

Answer 1

__constant

限定符用于恒定内存，有些卡在纹理缓存中获取它并从 __global 获得独立带宽，但大小非常有限。

__global __read_only * float

意味着，如果硬件认为合适，opencl 实现将尝试将其放入缓存中（或使用其他数据路径），但它是 __global，因此仅受 vram 大小或其分数的限制，而 __constant 仅受 64kB（例如）限制。

这些限定符用于设备端优化。

在主机端优化时，您应该为它提供一个

CL_MEM_READ_ONLY 

作为缓冲区创建的标志。这意味着设备只能从中读取（可能使用一些 DMA/pcie 访问/缓存优化），但可以使用 enqueuewrite 或 map unmap 操作从主机端写入（作为 C# C++ 代码之类的主机，而不是设备）。

__constant

用于参数常量定义，而不用于要处理的数据。

如果您正在编写过滤器代码，则数据可能是 __global 并且过滤器掩码可能是 __constant 如果不能放入 __private 内存（具有最终带宽）或 __local 内存（比私有慢），因此访问掩码字节不会减少数据带宽。

现在回答您的问题：

" a_buffer 是全局内存还是常量内存？"

它对于设备端（内核端）是全局的，因为您将其声明为 __global 但它可以在主机端（硬件）的任何位置。

编辑：对于主机端，取决于使用了哪些其他标志，例如，USE_HOST_PTR 使其可以从系统 RAM 直接访问，并且设备端只有一个虚拟缓冲区，没有它和只有一个 CL_MEM_READ_WRITE 设备内存将有一个真正的缓冲区及其在 RAM 中的映射影子（作为 clenqueueread 或 clenqueuewrite 的子步骤），复制将首先访问这个影子，然后上传到 gpu。

示例设备：4GB DDR3L 笔记本电脑中的 Intel(R) HD (TM) GRAPHICS 400：

Query                                           value
CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE                 65536 bytes
CL_DEVICE_GLOBAL_MEM_CACHE_SIZE                   262144 bytes
CL_DEVICE_GLOBAL_MEM_SIZE                     1636414260 bytes

CL_DEVICE_GLOBAL_MEM_CACHE_TYPE               CL_READ_WRITE_CACHE
CL_DEVICE_LOCAL_MEM_SIZE                      65536(vs constant, benchmark it)
CL_DEVICE_LOCAL_MEM_TYPE                      CL_LOCAL(so is faster than global)  

您无法查询私有内存大小，但对于中段游戏 AMD 卡，它是每个线程组 256kB。如果您设置每组 64 个线程，它可以使用每个线程或一半的 4kB 寄存器空间（因为编译器优化），然后因为溢出到全局内存而变慢。