OpenCL：传递指向本地内存的指针

如何解决OpenCL：传递指向本地内存的指针

我有以下示例代码：

int compute_stuff(int *array)
{
    /* do stuff with array */
    ...
    return x;
}

__kernel void my_kernel()
{
    __local int local_mem_block[LENGTH*MY_LOCAL_WORK_SIZE];
    int result;

    /* do stuff with local memory block */
    result = compute_stuff(local_mem_block + (LENGTH*get_local_id(0)));
    ...
}

以上示例在我的 NVIDIA 卡 (RTX 2080) 上编译和执行良好。
但是当我尝试在带有 AMD 卡的 Macbook 上进行编译时，出现以下错误：

error: passing '__local int *' to parameter of type '__private int *' changes address space of pointer

好的，那么我将“compute_stuff”函数更改为以下内容：

int compute_stuff(__local int *array)

现在NVIDIA和AMD都编译好了，没问题... 但后来我又做了一个测试，在同一台 Macbook 上使用 WINE 编译它（而不是在 bootcamp 中引导到 Windows），它给出了以下错误：

error: parameter may not be qualified with an address space

所以看起来好像不应该用地址空间限定函数参数。很公平。但是，如果我不这样做，那么本机 Windows 上的 AMD 会认为我正在尝试将指针的地址空间更改为 private（我猜是因为它假定所有函数参数都是私有的？）。

什么是处理这个问题的好方法，以便所有三个环境都乐于编译它？作为最后的手段，我正在考虑简单地让程序检查构建是否在没有限定符的情况下失败，如果是，则替换“__local”限定符并再次构建......看起来像一个黑客，但它可以工作。

解决方法

int* 中的 int compute_stuff(int *array) 是 __generic 地址空间。调用 result = compute_stuff(local_mem_block+...); 将其隐式转换为 __local，这是根据 OpenCL 2.0 Khronos specification 允许的。

可能是 AMD 默认使用 OpenCL 1.2。可能在 –cl-std=CL2.0 或 clBuildProgram() 中明确设置 clCompileProgram()。

为了使代码与 OpenCL 1.2 兼容，您可以将函数中的指针显式设置为 __local：int compute_stuff(__local int *array)。 OpenCL 允许将函数参数设置到地址空间 __global 和 __local。 WINE 似乎在这里有一个错误。也许内联函数可以解决它：int __attribute__((always_inline)) compute_stuff(__local int *array).

作为最后的手段，您可以执行您建议的方法。您可以检测它是否在 WINE 系统上运行，例如 this。这样，您就可以在两种代码变体之间切换，而无需进行两次编译和检测错误。

,

我同意 ProjectPhysX 的观点，它似乎是 WINE 实现的一个错误。我还发现以下内容似乎满足所有三种环境：

int compute_stuff(__local int * __private array)
{
    ...
}

__kernel void my_kernel()
{
    __local int local_mem_block[LENGTH*MY_LOCAL_WORK_SIZE];
    __local int * __private samples;

    samples = local_mem_block + (LENGTH*get_local_id(0));

    result = compute_stuff(samples);
}

上面明确指出指针本身是private，而它指向的内存保存在local地址空间中。所以这消除了任何歧义。