NVIDIA __constant 内存：如何在 OpenCL 和 CUDA 中从主机填充常量内存？

Question

我在主机上有一个缓冲区（数组），它应该驻留在设备（在本例中为 NVIDIA GPU）的常量内存区域中。

那么，我有两个问题：

如何分配一块常量内存？考虑到我正在跟踪设备上的可用常量内存，并且我知道，事实上，我们有那么多可用内存（此时）
如何根据主机上运行时间计算的值初始化（填充）这些数组？

我在网上搜索了这个，但是没有简明的文档记录这个。如果提供的示例同时在 OpenCL 和 CUDA 中，我将不胜感激。 OpenCL 的示例对我来说比 CUDA 更重要。

Answer 1

How can I allocate a chunk of constant memory? Given the fact that I am tracing the available constant memory on the device and I know, for a fact, that we have that amount of memory available to us (at this time)

在 CUDA 中，你不能。没有常量内存的运行时分配，只有通过 __constant__ 说明符静态定义内存，这些说明符在汇编时映射到常量内存页面。您可以在运行时生成一些包含此类静态声明的代码并通过 nvrtc 编译它，但这似乎需要付出很多努力64kb。静态声明一个 64kb 常量缓冲区并在运行时根据需要使用它似乎更简单（至少对我而言）。

How can I initialize (populate) those arrays from values that are computed at the runtime on the host?

如评论中所述，请参阅。 cudaMemcpyToSymbol API 就是为此目的创建的，它的工作方式与标准 memcpy 一样。

在功能上，OpenCL 中的__constant 和CUDA 中的__constant__ 没有区别。同样的限制适用：编译时的静态定义（标准 OpenCL 执行模型中的运行时），64kb 限制。

Answer 2

免责声明：我无法在 CUDA 方面为您提供帮助。

对于 OpenCL，从 programmer/API 的角度来看，constant 内存实际上被视为 read-only global 内存，或者在内核源代码中内联定义。

在内核代码中定义常量变量、数组等，如constant float DCT_C4 = 0.707106781f;。请注意，如果您愿意，可以在运行时在主机上动态生成内核代码以生成派生常量数据。
通过缓冲区对象将常量内存从主机传递到内核，就像您对 global 内存所做的那样。只需在内核函数原型的 constant 内存区域中指定一个指针参数，并在主机端使用 clSetKernelArg() 设置缓冲区，例如：

kernel void mykernel(
    constant float* fixed_parameters,
    global const uint* dynamic_input_data,
    global uint* restrict output_data)
{

    cl_mem fixed_parameter_buffer = clCreateBuffer(
        cl_context,
        CL_MEM_READ_ONLY | CL_MEM_HOST_NO_ACCESS | CL_MEM_COPY_HOST_PTR,
        sizeof(cl_float) * num_fixed_parameters, fixed_parameter_data,
        NULL);
    clSetKernelArg(mykernel, 0, sizeof(cl_mem), &fixed_parameter_buffer);

请确保考虑到 CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE 所报告的所用上下文的值！使用 constant 内存缓冲区来传输输入数据通常没有帮助，最好将其存储在 global 缓冲区中，即使它们被标记为 read-only 用于内核。 constant 内存对于大部分 work-items 使用的数据最有用。通常有一个相当严格的大小限制，例如 64KiB - 如果你试图超过这个，一些实现可能会“溢出”到全局内存，这将失去你使用常量内存获得的任何性能优势。

Answer 3

对于 cuda，我使用 驱动程序 API 和 NVRTC 并使用这样的全局常量数组创建内核字符串：

auto kernel = R"(
..
__constant__ @@Type@@ buffer[@@SIZE@@]={
   @@elm@@
};
..
__global__ void test(int * input)
{   }

)";

然后在run-time和compile中用大小和元素值信息替换@@-pattern单词，像这样：

__constant__ int buffer[16384]={ 1,2,3,4, ....., 16384 };

所以主机是run-time，设备是compile-time。缺点是内核字符串太大，可读性差，连接类需要显式链接（就像你正在编译一个 C++ 项目一样）其他编译单元。但是对于只有自己实现的简单计算（没有直接使用host-definitions），它与运行时相同API.

由于大字符串需要额外的解析时间，您可以缓存ptx中间数据，也可以缓存ptx生成的二进制文件。然后您可以检查内核字符串是否已更改并且需要 re-compiled.

您确定 __constant__ 值得付出努力吗？您是否有一些基准测试结果表明它确实提高了性能？（过早的优化是万恶之源）。也许您的算法适用于 register-tiling 而数据来源无关紧要？

NVIDIA __constant 内存：如何在 OpenCL 和 CUDA 中从主机填充常量内存？

NVIDIA __constant memory: how to populate constant memory from host in both OpenCL and CUDA?

cuda

gpu

nvidia

opencl

gpu-constant-memory