如何在 OpenCL 中声明本地内存？答案

【问题标题】：How to declare local memory in OpenCL?如何在 OpenCL 中声明本地内存？
【发布时间】：2012-02-11 21:47:50
【问题描述】：

我正在运行下面的 OpenCL 内核，二维全局工作大小为 1000000 x 100，本地工作大小为 1 x 100。

__kernel void myKernel(
        const int length, 
        const int height, 
        and a bunch of other parameters) {

    //declare some local arrays to be shared by all 100 work item in this group
    __local float LP [length];
    __local float LT [height];
    __local int bitErrors = 0;
    __local bool failed = false;

    //here come my actual computations which utilize the space in LP and LT
}

然而这拒绝编译，因为参数length 和height 在编译时是未知的。但是我完全不清楚如何正确地做到这一点。我应该在 memalloc 中使用指针吗？如何以只为整个工作组分配一次内存而不是每个工作项一次的方式来处理这个问题？

我只需要 2 个浮点数组，1 个整数和 1 个布尔值，它们在整个工作组之间共享（所以所有 100 个工作项）。但是我找不到任何可以正确执行此操作的方法...

【问题讨论】：

stackoverflow.com/questions/2541929/…

标签： memory opencl

【解决方案1】：

这比较简单，你可以将本地数组作为参数传递给你的内核：

kernel void myKernel(const int length, const int height, local float* LP, 
                     local float* LT, a bunch of other parameters)

然后，您使用NULL 的value 和size 设置内核参数，该size 等于您要为参数分配的大小（以字节为单位）。因此应该是：

clSetKernelArg(kernel, 2, length * sizeof(cl_float), NULL);
clSetKernelArg(kernel, 3, height* sizeof(cl_float), NULL);

本地内存总是由工作组共享（而不是私有），所以我认为bool 和int 应该没问题，但如果不是，您也可以始终将它们作为参数传递。

与您的问题并不真正相关（也不一定相关，因为我不知道您打算在什么硬件上运行它），但至少 gpus 并不特别喜欢不是特定幂的倍数的工作大小两个（我认为 nvidia 是 32，amd 是 64），这意味着这可能会创建包含 128 个项目的工作组，其中最后 28 个基本上被浪费了。因此，如果您在 gpu 上运行 opencl，如果您直接使用大小为 128 的工作组（并适当更改全局工作大小），可能会提高性能

附带说明：我一直不明白为什么每个人都使用下划线变体来表示kernel, local and global，这对我来说似乎更难看。

【讨论】：

嗯...那么我应该如何解决这个工作组大小问题？ 100 是一个标称值，它可能会根据问题的具体实例而有所不同，但我需要全局输入的每个 1x100 子块的这些本地内存变量。我假设如果该子块不是工作组，则无法在其正确的 1x100 子块之间共享变量？（至于旁注，我从来没有尝试过没有__，谢谢你的提示！）
我可能应该将二维数组传递给内存，但我该怎么做呢？我可以按照clSetKernelArg(kernel, 2, length * local_work_size[0] * sizeof(cl_float), NULL); 的方式传递一些东西，但这是一个一维数组。将其设为二维可能会更好吗？
另外，bool 和 int 也不好，我得到一个 error: variable "bitErrors" may not be initialized。如果我用上面的本地语句替换它，我会得到error: a parameter cannot be allocated in a named address space。当然，我可以将其设为长度为 1 的数组，但这并不是最漂亮的解决方案... :-)
@user1111929：您说得对，您只能在工作组内共享变量。我刚刚提到它，以防你可以改变你的工作大小。并且将本地数组作为一维传递确实是要走的路，opencl 并不真正支持嵌套指针。关于bool 和int：我的意思是如果它不能像那样工作，就将它作为指针传递。
我相信第二个clSetKernelArg() 调用的第二个参数应该是3。正如所发布的，这将设置相同的参数两次。

【解决方案2】：

你也可以这样声明你的数组：

__local float LP[LENGTH];

并在内核编译中将 LENGTH 作为定义传递。

int lp_size = 128; // this is an example; could be dynamically calculated
char compileArgs[64];
sprintf(compileArgs, "-DLENGTH=%d", lp_size);
clBuildProgram(program, 0, NULL, compileArgs, NULL, NULL);

【讨论】：

了解这一点很有用，但没有解决有关如何在运行时（而不是编译时）动态分配的问题。
@MasterHD OpenCL 代码可以（并且确实是最简单的示例）在运行时编译，所以这个答案确实解决了这个问题。

【解决方案3】：

您不必在内核之外分配所有本地内存，尤其是当它是一个简单变量而不是数组时。

您的代码无法编译的原因是 OpenCL 不支持本地内存初始化。这是在文档中指定的 (https://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/local.html)。在CUDA中也不可行(Is there a way of setting default value for shared memory array?)

ps：Grizzly 的回答已经足够好了，如果我能把它作为评论发布就更好了，但是我受到声誉政策的限制。对不起。

【讨论】：