cuda - 在共享内存中创建数组，无需模板，如 PyOpenCL

Question

如何在不修改内核的情况下使用官方示例中看到的模板在共享内存中创建数组。还是以官方方式使用模板？

在 PyOpenCL 中，我可以通过设置内核参数在本地内存中创建一个数组

kernel.set_arg(1,numpy.uint32(a_width))

... 
KERNEL_CODE = """
__kernel void matrixMul(__local float* A_temp,...)
    { ...} """

Answer 1

CUDA 在内核运行时支持动态共享内存分配，但机制与 OpenCL 有点不同。在 CUDA 运行时 API 中，使用动态分配/大小的共享内存和启动来调整内存大小的内核使用以下语法：

__global__ void kernel(...)
{
    extern __shared__ typename buffer[];

    ....
}
....
kernel <<< griddim, blockdim, sharedmem, streamID >>> (...)

其中sharedmem是将分配给缓冲区的每个块的总字节数。

在 PyCUDA 中，相同的机制是这样工作的：

mod = SourceModule("""
    __global__ void kernel(...)
    {
        extern __shared__ typename buffer[];

        ....
    }
  """)

func = mod.get_function("kernel")
func.prepare(..., shared=sharedmem)
func.prepared_call(griddim,blockdim,...)

将共享内存分配大小传递给prepare方法。

Answer 2

我不完全理解这个问题。我不使用 Python，但非常了解 OpenCL。

在 OpenCL 中，您有两种创建共享/本地内存缓冲区的可能性：

1）您在问题中添加了一个内核参数。2）在内核本身内静态定义一个缓冲区，例如：

__local buffer[1024];

使用 OpenCL 没有其他机会可以做到这一点。如何创建内核代码字符串以将其传递给 OpenCL 是另一个与 Python 相关的问题。我不是这方面的专家。