如何使用 PyOpenCL 将带有数组和变量的 C 结构传递给 OpenCL 内核

Posted 2023-02-23

【中文标题】如何使用 PyOpenCL 将带有数组和变量的 C 结构传递给 OpenCL 内核

【英文标题】：How to pass a C struct with arrays and variables to OpenCL kernel using PyOpenCL

【发布时间】：2018-06-17 03:05:41

【问题描述】：

因此，我必须使用 PyOpenCL 或使用 Python 的一些解决方法将一些数据传递给 OpenCL 内核。数据在内核端作为结构读取，我无法更改内核，因为它工作正常，并且是我的代码必须使用的更大项目的一部分。

内核是这样的：

typedef struct VglClStrEl 
    float data[VGL_ARR_CLSTREL_SIZE];
    int ndim;
    int shape[VGL_ARR_SHAPE_SIZE];
    int offset[VGL_ARR_SHAPE_SIZE];
    int size;
 VglClStrEl;

typedef struct VglClShape 
    int ndim;
    int shape[VGL_ARR_SHAPE_SIZE];
    int offset[VGL_ARR_SHAPE_SIZE];
    int size;
 VglClShape;

__kernel void kernel(__global unsigned char* img_input, 
                     __global unsigned char* img_output,  
                     __constant VglClShape* img_shape,
                     __constant VglClStrEl* window)


    // do what is needed

因此，如您所见，VglClShape 和 VglClStrEl 结构具有不同的类型数组和静态位大小变量。

[1] 解决方法支持只有一种类型数组的结构（或者我很遗憾地未能找到一种方法来处理多种数组类型）。

[2] 解决方法是有关如何将 Python 数据传递到 OpenCL 内核结构的 PyOpenCL 文档参考。这种方法根本不支持数组。

那么，如何传递 OpenCL 内核可以读取的 Python 数据？我已经拥有 Python 端的所有数据，我只需要知道如何将其从 Python 传递到内核。

在您问之前：我正在使用 Python 3，并且我无法更改内核。

是的，数组大小是静态的。你可以假设这样的事情：

VGL_ARR_CLSTREL_SIZE=256;
VGL_ARR_SHAPE_SIZE=20;

[1]Passing struct with pointer members to OpenCL kernel using PyOpenCL

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

【问题讨论】：

可以使用标准的ctypes模块吗？

@PM2Ring 众所周知，ctypes 是一个 Python 外部库，用于在 Python 上使用 C 语言。我需要做的是相反的事情：以它理解的方式将 Python 数据传递给 OpenCL（注意两者之间的内存分配对称性）。我可以使用 ctypes 来做类似的事情吗？

我不了解 OpenCL，但我使用 ctypes 与 OpenSSL 库进行交互。有关示例，请参阅here。而最终的代码块here还有一个例子。

【参考方案1】：

有一种骇人听闻的方法可以做到这一点，这需要一些繁琐的字节争论。大概您可以部署一个小型 OpenCL 探测内核？ （在任何情况下，PyOpenCL 都会在某些操作的后台执行此操作）

基本思路是：

了解 OpenCL 设备如何通过运行单个实例内核来对齐结构的所有元素 创建一个 numpy 字节数组以匹配 OpenCL 结构的大小 按字节将 Python 结构的每个元素复制到此数组中 在调用不可更改的 OpenCL 内核时，通过一袋字节缓冲区传递此数组

以下内核可以完成这项工作：

__kernel void get_struct_sizes( __global uint *struct_sizes )

    const uint global_id = get_global_id(0u)+get_global_id(1u)*get_global_size(0u);
    VglClStrEl vgclstrel;
    VglClShape vgclshape;
    uint offset;

    printf("In GPU (probing):\n Kernel instance = %d\n", global_id);

    if (global_id==0) 
        offset = (uint)&(vgclstrel.data);
        struct_sizes[0] = (uint)sizeof(vgclstrel);
        struct_sizes[1] = (uint)&(vgclstrel.ndim)-offset;
        struct_sizes[2] = (uint)&(vgclstrel.shape)-offset;
        struct_sizes[3] = (uint)&(vgclstrel.offset)-offset;
        struct_sizes[4] = (uint)&(vgclstrel.size)-offset;
        offset = (uint)&(vgclshape.ndim);
        struct_sizes[5] = (uint)sizeof(vgclshape);
        struct_sizes[6] = (uint)&(vgclshape.shape)-offset;
        struct_sizes[7] = (uint)&(vgclshape.offset)-offset;
        struct_sizes[8] = (uint)&(vgclshape.size)-offset;
    
    return;

执行这个内核并将struct_sizes返回到vgclshape_sizes，创建这个数组：

img_shape  = np.zeros((vgclshape_sizes[0]), dtype=np.uint8)

然后把你需要的复制进去：

def copy_into_byte_array(value, byte_array, offset):
        for i,b in enumerate(np.ndarray.tobytes(value)):
            byte_array[i+offset] = b
copy_into_byte_array(ndim,   img_shape, 0) 
copy_into_byte_array(shape,  img_shape, vgclshape_sizes[1]) 
copy_into_byte_array(offset, img_shape, vgclshape_sizes[2]) 
copy_into_byte_array(size,   img_shape, vgclshape_sizes[3]) 

我在这里跳过了一些步骤；填写它们，您会发现这种方法有效。我能够将演示结构传递给您不可侵犯的内核的虚拟副本。

我很想知道是否有更优雅的方法来执行任何/所有这些步骤。我还希望字节顺序等方面会出现问题，否则这些问题将是透明的。运气好的话，你可以解决它们。

【讨论】：

你是如何将字节包传递给内核的？当我尝试传递字节包时，我收到错误 clSetKernelArg failed: INVALID_ARG_SIZE

看这里：github.com/cstarknyc/PyPlayground/tree/master/Pass_Struct ...但基本上：struct_sizes = np.zeros(9, dtype=np.uint32)struct_sizes_buffer = cl.Buffer(cl_state.context, COPY_WRITE, hostbuf=struct_sizes) buffer_list = [struct_sizes_buffer] cl_state.kernel.set_args(*buffer_list) cl_state.kernel.set_scalar_arg_dtypes( [None]*len(buffer_list) ) cl_state.event = cl.enqueue_nd_range_kernel(cl_state.queue, cl_state.kernel, global_size, local_size)

终于让它按预期工作了。非常感谢您对 Colin 的帮助！

还不行。我没有15点声望。当我得到它们时，我会投票给你