如何将二维数组传递到pycuda中的内核？

Question

我找到了答案here，但目前尚不清楚我是否应该重塑阵列。在将2d数组传递给pycuda内核之前，是否需要将2d数组重塑为1d？

Answer 1

没有必要重塑2D gpuarray以将其传递给CUDA内核。

正如我在你所链接的答案中所说，2D numpy或PyCUDA数组只是一个音调线性内存的分配，默认情况下以行主要顺序存储。两个都有两个成员告诉你访问数组所需的一切 - shape和strides。例如：

In [8]: X=np.arange(0,15).reshape((5,3))

In [9]: print X.shape
(5, 3)

In [10]: print X.strides
(12, 4)

形状是自我解释的，步幅是以字节为单位的存储间距。内核代码的最佳实践是将PyCUDA提供的指针视为使用cudaMallocPitch分配，并将stride的第一个元素视为内存中行的字节间距。一个简单的例子可能如下所示：

import pycuda.driver as drv
from pycuda.compiler import SourceModule
import pycuda.autoinit
import numpy as np

mod = SourceModule("""
__global__ void diag_kernel(float *dest, int stride, int N)
{
    const int tid = threadIdx.x + blockDim.x * blockIdx.x;

    if (tid < N) {
    float* p = (float*)((char*)dest + tid*stride) + tid;
        *p = 1.0f;
    }
}
""")

diag_kernel = mod.get_function("diag_kernel")

a = np.zeros((10,10), dtype=np.float32)
a_N = np.int32(a.shape[0])
a_stride = np.int32(a.strides[0])
a_bytes = a.size * a.dtype.itemsize
a_gpu = drv.mem_alloc(a_bytes)
drv.memcpy_htod(a_gpu, a)
diag_kernel(a_gpu, a_stride, a_N, block=(32,1,1))
drv.memcpy_dtoh(a, a_gpu)

print a

这里有一些内存分配在设备上，一个归零的2D数组直接复制到该分配，内核的结果（用1填充对角线）复制回主机并打印。没有必要在过程中的任何点处展平或以其他方式修改2D numpy数据的形状或存储器布局。结果是：

$ cuda-memcheck python ./gpuarray.py 
========= CUDA-MEMCHECK
[[ 1.  0.  0.  0.  0.  0.  0.  0.  0.  0.]
 [ 0.  1.  0.  0.  0.  0.  0.  0.  0.  0.]
 [ 0.  0.  1.  0.  0.  0.  0.  0.  0.  0.]
 [ 0.  0.  0.  1.  0.  0.  0.  0.  0.  0.]
 [ 0.  0.  0.  0.  1.  0.  0.  0.  0.  0.]
 [ 0.  0.  0.  0.  0.  1.  0.  0.  0.  0.]
 [ 0.  0.  0.  0.  0.  0.  1.  0.  0.  0.]
 [ 0.  0.  0.  0.  0.  0.  0.  1.  0.  0.]
 [ 0.  0.  0.  0.  0.  0.  0.  0.  1.  0.]
 [ 0.  0.  0.  0.  0.  0.  0.  0.  0.  1.]]
========= ERROR SUMMARY: 0 errors