如何在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) 

然后,您可以用该kernelargument value的NULL和size等于要分配的参数(以字节)的大小.因此它应该是:

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

本地内存总是由工作组共享(而不是私有),所以我认为bool并且int应该没问题,但如果没有,你也总是可以将它们作为参数传递.

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

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