将向量（float4）内核参数传递给OpenCL（Python）

Question

有没有简单的方法可以把 float4 或其他向量参数传递给 OpenCL 内核？对于标量参数（比如 int、float），你可以在调用内核时直接传递。而对于数组参数，你需要先用 cl.Buffer() 把它复制到 GPU，然后再传递指针。当然，可能也可以像传数组那样传 float4，但我想知道有没有更简单、更清晰的方法（特别是使用 Python、numpy 和 pyOpenCL 时）。

我尝试把大小为 4 的 numpy 数组作为 float4 传递，但没有成功。有没有其他方法可以做到这一点？

例如：

内核：

__kernel void myKernel( __global float  * myArray, float myFloat, float4 myFloat4 )

Python：

myFloat4   = numpy.array  ( [1.0 ,2.0 ,3.0], dtype=np.float32 ) 
myArray    = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=myArray_host)
kernelargs = ( myArray , numpy.float32(myFloat) , myFloat4) 
prg.myKernel(queue, cl_myArray.shape() , None, *(kernelargs) )

我遇到了错误：

pyopencl.LogicError: when processing argument #2 (1-based): clSetKernelArg failed: invalid arg size

另一种可能性是把它作为一组标量 int 或 float 传递，比如：

__kernel void myKernel( __global float  * myArray, float myFloat, float myFloat4_x, float myFloat4_y, float myFloat4_z  )

kernelargs = ( myArray , numpy.float32(myFloat) ,numpy.float32(myFloat4_x),numpy.float32(myFloat4_y),numpy.float32(myFloat4_z))

但这样也不太方便——如果你想传递 4 个 float4 和 5 个 int3 给内核，你很容易就会在很多变量名中迷失。

我觉得在 OpenCL 中传递 int 和 float 的向量（2、3、4）应该是相当常见的，比如 3D 数据网格的大小。所以我在想，真的有必要通过 cl.Buffer() 作为指针来传递吗？

我猜常量参数 float4 的速度也比 *float 快，因为它可以被所有工作项共享作为常量。

Answer 1

我注意到三件事：

1. 从错误信息来看，第二个内核参数，也就是 myFloat 似乎有问题。如果你在内核签名中把它声明为 const 参数，会发生什么？如果你这样做，会有什么变化呢？

   myFloat = myFloat.astype(np.float32)
   kernelArgs = (..., myFloat, ...)
   prg.myKernel(...)
   

2. 你想定义一个包含四个元素的向量 myFloat4，但你只给了三个值 [1.0, 2.0, 3.0]。另外，试着在内核签名中设置 const float4 myFloat4。

3. 在实际调用内核时，kernelargs 元组不需要额外的括号：

   prg.myKernel(queue, cl_myArray.shape() , None, *kernelargs)
   

Answer 2

问题出在这里：

myFloat4   = numpy.array  ( [1.0 ,2.0 ,3.0], dtype=numpy.float32 )

但是我的 myFloat4.size 等于 3

只需输入这个：

myFloat4   = numpy.array  ( [1.0 ,2.0 ,3.0, 4.0], dtype=numpy.float32 )

其余的代码都没问题

Answer 3

我觉得在Python中创建一个float4（四个浮点数的组合）是个不错的方法：

import numpy as np
import pyopencl as cl
import pyopencl.array as cl_array

data= np.zeros(N, dtype=cl_array.vec.float4)

补充：为了提供一个最小可工作示例（MWE）：

import numpy as np
import pyopencl as cl
import pyopencl.array as cl_array


deviceID = 0
platformID = 0
workGroup=(1,1)

N = 10
testData = np.zeros(N, dtype=cl_array.vec.float4)

dev = cl.get_platforms()[platformID].get_devices()[deviceID]

ctx = cl.Context([dev])
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags
Data_In = cl.Buffer(ctx, mf.READ_WRITE, testData.nbytes)


prg = cl.Program(ctx, """

__kernel void   Pack_Cmplx( __global float4* Data_In, int  N)
{
  int gid = get_global_id(0);

  Data_In[gid] = 1;
}
 """).build()

prg.Pack_Cmplx(queue, (N,1), workGroup, Data_In, np.int32(N))
cl.enqueue_copy(queue, testData, Data_In)


print testData