OpenCL：将指针传递给本地内存答案

【问题标题】：OpenCL: Passing a pointer to local memoryOpenCL：将指针传递给本地内存
【发布时间】：2021-07-12 01:15:05
【问题描述】：

我有以下示例代码：

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 都可以正常编译了，没问题... 但后来我又进行了一项测试，使用 WINE 在同一台 Macbook 上编译它（而不是在 bootcamp 中启动到 Windows），它给出了以下错误：

error: parameter may not be qualified with an address space

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

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

【问题讨论】：

标签： pointers memory opencl gpgpu

【解决方案1】：

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

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

为了保持代码与 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。这样，您就可以在两个代码变体之间切换，而无需编译两次并检测到错误。

【讨论】：

【解决方案2】：

我同意 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 地址空间中。所以这消除了任何歧义。

【讨论】：