opencl - Nvidia 上 opencl 共享内存中的动态分配

Question

我正在按照示例 here创建一个可变长度的本地内存数组。内核签名是这样的:

__kernel void foo(__global float4* ex_buffer,
                 int ex_int,
                 __local void *local_var) 

然后我为本地内存内核参数调用 clSetKernelArg，如下所示:

clSetKernelArg(*kern, 2, sizeof(char) * MaxSharedMem, NULL)

MaxSharedMem 是通过查询 CL_DEVICE_LOCAL_MEM_SIZE 设置的。然后在内核中，我将分配的本地内存分成几个数组和其他数据结构，并在我认为合适的时候使用它们。所有这些都适用于 AMD(gpu 和 cpu)和 Intel 设备。但是，在 Nvidia 上，当我将该内核入队然后在队列上运行 clFinish 时，我收到错误 CL_INVALID_COMMAND_QUEUE。

这是一个生成上述错误的简单内核(本地工作大小为 32):

__kernel 
void s_Kernel(const unsigned int N, __local void *shared_mem_block ) 
{
    const ushort thread_id = get_local_id(0);   
    __local double *foo = shared_mem_block;
    __local ushort *bar = (__local ushort *) &(foo[1000]);


    foo[thread_id] = 0.;
    bar[thread_id] = 0;
}

如果我在本地内存中静态分配相同的数组和数据结构，内核运行良好。有人可以为此行为和/或解决方法提供解释吗？

Answer 1

对于那些感兴趣的人，我终于收到了 Nvidia 的解释。当通过 void 指针传入共享内存块时，实际对齐与指向 double 指针(8 字节对齐)的预期对齐不匹配。由于未对齐，GPU 设备抛出异常。

正如其中一条评论所指出的，避免该问题的一种方法是让内核参数成为指向编译器将正确对齐到至少 8 个字节(double、ulong 等)的内容的指针。

理想情况下，编译器会负责任何特定于设备的对齐问题，但由于在我的问题中提到的小内核中有一个隐式指针转换，我认为它会造成混淆。

一旦内存按 8 字节对齐，转换为采用较短对齐方式(例如 ushort)的指针类型就不会出现问题。因此，如果您像我一样链接内存分配，并且指针指向不同的类型，请确保指针指向内核签名中的最大类型。