OpenCL：使用什么类型的内存？答案

【问题标题】：OpenCL: What type of memory to use?OpenCL：使用什么类型的内存？
【发布时间】：2015-02-16 18:34:17
【问题描述】：

我有某种过滤内核，像这样：

__kernel void filterKernel (__global float4 *filter, __global float4* in_array, __global float4* out_array)
{
...
out_array[tid] = in_array[tid] * filter[fid];
...
}

kernel filterKernel被调用多次（约1000次）。
变量 filter 是一个浮点数组，从不改变它的值（对于所有工作组和所有内核调用都保持相同）。
in_array 包含 32768 个浮点数。

这个变量filter的最佳声明是什么？ __持续的？ __当地的？也许在这里和那里放置一个“const”？什么对编译器帮助最大？是什么让代码最快？

【问题讨论】：

您的问题得到解答了吗？如果是，请接受答案。

标签： memory constants local opencl

【解决方案1】：

您应该使用常量地址空间 (__constant)，因为大多数 GPU 都有用于常量内存的特殊缓存。唯一的问题是常量内存很小（大约 16-64KB）。

【讨论】：

如果我使用常量内存：我应该在我的 .cl 文件中添加一个定义（例如 ___constant float4* filter = {1.0, 2.0, 3.0, ...}; ），还是保留 filter 在我的函数接口中并在调用函数之前填充 CPU 上的数组？有区别吗？
@Frizz：使用 clSetKernelArg 从 CPU 填充它。
这比在 .cl 文件中填充 __constant 数组要慢得多。请参阅下面的我的 cmets。

【解决方案2】：

__local 将是错误的，因为您无法将其初始化为任何内容。您可能想使用 __constant，前提是它适合。

【讨论】：

【解决方案3】：

如果不是太大，请尝试在 .cl 文件中全局定义您的 过滤器。
在那里，您可以尝试在 __constant 或 __local 空间中分配它，并比较哪个更快。但并不是所有的 SDK 都支持 __local 地址空间中的全局变量（我在看你的 ATI）。

如果您仍想将过滤器作为内核参数传递，请考虑仅调用它的 SetKernelArg(0, ...) 一次。只要内核参数的值或索引不变，也没有必要调用 SetKernelArg() 1000 次。虽然这可能不会对性能产生可衡量的影响，但它仍然更干净。

【讨论】：

到目前为止，我尝试过的所有各种方法（常量内存、LDS、纹理缓存）在 .cl 文件中将其定义为 __constant，是迄今为止（！）最快的。令人惊讶的是，A）在调用它之前通过 SetkernelArg() 将 __constant 数组传递给内核 - 或 B）在 .cl 文件本身中定义 __constant 数组之间存在巨大（！）差异。不幸的是，方法 B) 仅适用于 ATI/AMD。 Nvidia 的 OpenCL 实现在这里有错误（我已经报告了这个错误，将在 CUDA 3.3 中修复）。
是的，我同意这一点。 CUDA SDK 有很多错误。我尝试使用位移操作在内核中实现 byte_swap() 函数。 OpenCL 编译器识别该函数的用途并尝试将其映射到 bswap32 resp。 bswap64 指令 - 无法找到或解决...
Frizz：您是否检查过设置 __constant 缓冲区的内核参数仅一次与在每个内核入队之前设置的影响？
不，我没有检查这是否有影响。会做。 ... 通常我只测试 clEnqueueNDRange() 调用完成的时间。
已检查：完全没有区别（仅在每次内核调用之前设置 __constant 缓冲区内核参数一次）。至少在我的硬件（GT240）上没有。