使用 PyOpenCL 将带有指针成员的结构传递给 OpenCL 内核

Question

假设我有一个内核来计算两个数组的元素总和。我没有将 a、b 和 c 作为三个参数传递，而是将它们设为结构成员，如下所示：

typedef struct
{
    __global uint *a;
    __global uint *b;
    __global uint *c;
} SumParameters;

__kernel void compute_sum(__global SumParameters *params)
{
    uint id = get_global_id(0);
    params->c[id] = params->a[id] + params->b[id];
    return;
}

如果您使用 PyOpenCL [1] 的 RTFM，则有关于结构的信息，其他人也已经解决了这个问题 [2] [3] [4]。但是我找到的 OpenCL 结构示例都没有指针作为成员。

具体来说，我担心主机/设备地址空间是否匹配，以及主机/设备指针大小是否匹配。有人知道答案吗？

[1]http://documen.tician.de/pyopencl/howto.html#how-to-use-struct-types-with-pyopencl

[2]Struct Alignment with PyOpenCL

[3]http://enja.org/2011/03/30/adventures-in-opencl-part-3-constant-memory-structs/

[4]http://acooke.org/cute/Somesimple0.html

Answer 1

不，不保证地址空间匹配。对于基本类型（float，int，...），您有对齐要求（标准的第 6.1.5 节），并且您必须使用 OpenCL 实现的 cl_type 名称（在 C 中编程时，pyopencl 完成了我的工作会说）。

对于指针，由于这种不匹配，它甚至更简单。标准 v 1.2 的第 6.9 节的开头（1.1 版的第 6.8 节）指出：

在程序中声明的内核函数的参数是指针 必须使用 __global、__constant 或 __local 限定符声明。

在 p 点：

声明为结构的内核函数的参数或 union 不允许 OpenCL 对象作为 结构或联合。

还要注意点d。：

具有灵活（或未调整大小）的可变长度数组和结构 不支持数组。

所以，没有办法让你的内核按照你的问题中描述的那样运行，这就是为什么你找不到一些 OpenCl 结构的例子有指针作为成员。
我仍然可以提出一个解决方法，利用内核在 JIT 中编译的事实。它仍然要求您正确打包数据并注意对齐，最后在程序执行期间大小不会改变。老实说，我会选择一个以 3 个缓冲区作为参数的内核，但无论如何，它就在那里。

这个想法是使用预处理器选项-D，如下面的python示例：

内核：

typedef struct {
    uint a[SIZE];
    uint b[SIZE];
    uint c[SIZE];
} SumParameters;

kernel void foo(global SumParameters *params){
    int idx = get_global_id(0);
    params->c[idx] = params->a[idx] + params->b[idx];
}

主机代码：

import numpy as np
import pyopencl as cl

def bar():
   mf = cl.mem_flags
   ctx = cl.create_some_context()
   queue = cl.CommandQueue(self.ctx)
   prog_f = open('kernels.cl', 'r')
   #a = (1, 2, 3), b = (4, 5, 6)          
   ary = np.array([(1, 2, 3), (4, 5, 6), (0, 0, 0)], dtype='uint32, uint32, uint32')
   cl_ary = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=ary)
   #Here should compute the size, but hardcoded for the example
   size = 3
   #The important part follows using -D option
   prog = cl.Program(ctx, prog_f.read()).build(options="-D SIZE={0}".format(size))    
   prog.foo(queue, (size,), None, cl_ary)
   result = np.zeros_like(ary)
   cl.enqueue_copy(queue, result, cl_ary).wait()
   print result

结果：

[(1L, 2L, 3L) (4L, 5L, 6L) (5L, 7L, 9L)]

Answer 2

我不知道自己的问题的答案，但我可以想到 3 种解决方法。我认为解决方法 3 是最佳选择。

解决方法 1：我们这里只有 3 个参数，所以我们可以只制作 a、b 和 c 内核参数。但是我读过你可以传递给内核的参数数量是有限制的，我个人喜欢重构任何需要超过 3-4 个参数的函数来使用结构（或者，在 Python 中，元组或关键字参数） .因此，此解决方案使代码更难阅读，并且无法扩展。

解决方法 2：将所有内容转储到一个巨大的数组中。然后内核看起来像这样：

typedef struct
{
    uint ai;
    uint bi;
    uint ci;
} SumParameters;

__kernel void compute_sum(__global SumParameters *params, uint *data)
{
    uint id = get_global_id(0);
    data[params->ci + id] = data[params->ai + id] + data[params->bi + id];
    return;
}

换句话说，不是使用指针，而是使用单个数组的偏移量。这看起来很像开始实现我自己的内存模型，感觉就像是在重新发明一个存在于 PyOpenCL 或 OpenCL 或两者中某处的轮子。

解决方法 3：制作 setter 内核。像这样：

__kernel void set_a(__global SumParameters *params, __global uint *a)
{
    params->a = a;
    return;
}

set_b、set_c 也是如此。然后使用 worksize 1 执行这些内核来设置数据结构。您仍然需要知道为参数分配多大的块，但如果它太大，不会发生任何不好的事情（除了一点点浪费的内存），所以我想说假设指针是 64 位。

这种解决方法的性能可能很糟糕（我想内核调用会产生巨大的开销），但幸运的是，这对我的应用程序来说并不重要（我的内核将一次运行几秒钟，这不是图形问题它必须以 30-60 fps 的速度运行，所以我想额外的内核调用设置参数所花费的时间最终将只是我工作量的一小部分，无论每个内核调用的开销有多高）。