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

Question

我正在运行下面的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个int和1个布尔值，它们在整个工作组之间共享（所以所有100个工作项）。但我没有找到任何正确做到这一点的方法......

Answer 1

它相对简单，您可以将本地数组作为参数传递给您的内核：

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

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

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

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

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

作为旁注：我从未理解为什么每个人都使用kernel, local and global的下划线变体，对我来说似乎更加丑陋。

Answer 2

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

您的代码无法编译的原因是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的答案已经足够好了，如果我可以将其作为评论发布，那会更好，但我受到声誉政策的限制。抱歉。

Answer 3

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

__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);