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

Question

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

19

我使用以下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 来进行指针操作吗？如何处理才能使内存只在整个工作组中分配一次，而不是每个工作项都分配一次？

我需要的仅仅是两个浮点数数组、1 个整数和 1 个布尔值，它们在整个工作组（即所有 100 个工作项）之间共享。但我未能找到任何正确处理这种情况的方法...

- user1111929

https://dev59.com/G3E85IYBdhLWcg3w64MA - Ciro Santilli OurBigBook.com

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

- David B

1

这是有用的知识，但并没有回答关于如何在运行时动态分配内存（而不是在编译时）的问题。 - MasterHD

1

@MasterHD OpenCL代码可以（而且大多数简单的示例确实是）在运行时编译，因此这个答案确实回答了这个问题。 - Dinei

1

当变量不是数组时，你无需在内核外分配所有本地内存。

你的代码无法编译的原因是 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的回答已经足够好了，如果我能够将其发布为评论，那将更好，但由于声望政策的限制，我受到了限制。抱歉。

- youwei

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Grizzly · Accepted Answer

这相对很简单，你可以将本地数组作为参数传递到您的内核中：

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, 3, height* sizeof(cl_float), NULL);

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

这跟您的问题并不相关（并且不一定相关，因为我不知道您计划在哪种硬件上运行），但至少显卡不喜欢工作大小不是某个2的幂次方倍数的情况（我想nvidia是32，amd是64），这意味着将会创建包含128个项目的工作组，其中最后28个项目基本上是浪费的。因此，如果您在GPU上运行opencl，直接使用大小为128的工作组（并相应更改全局工作大小）可能有助于提高性能。

顺便说一句：我从来没有理解为什么每个人都使用下划线变量写法表示kernel、local和global，对我来说看起来更丑陋。