OpenCL——Buffer和全局内存的区别

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 and READ_WRITE）和内存限定符（global and constant）有什么联系？

Answer 1

__constant

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

__global __read_only * float

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

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

在主机端优化时，您应该为其提供

CL_MEM_READ_ONLY

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

__constant

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

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

现在回答你的问题：

" is a_buffer a global memory or constant memory? "

它在设备端（内核端）是全局的 因为你将它声明为 __global 但它可以在主机端（硬件）的任何地方。

编辑：对于主机端，取决于使用了哪些其他标志，例如，USE_HOST_PTR 使其可以从系统 RAM 直接访问并且只有一个设备端的虚拟缓冲区，没有它并且只有 CL_MEM_READ_WRITE 设备内存将有一个真实的缓冲区及其在 RAM 中的映射影子（作为 clenqueueread 或 clenqueueewrite 的子步骤）并且复制将首先访问这个影子然后上传到 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 寄存器 space 每个线程或一半（由于编译器优化），然后由于溢出到全局内存而变慢。

OpenCL——Buffer和全局内存的区别

OpenCL - The difference between Buffer and global memory

global

constants

opencl