教程#

入门#

在使用PycUDA之前，必须导入并初始化：

import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule

注意你没有 have 使用 pycuda.autoinit 如果需要，也可以手动执行初始化、上下文创建和清理。

传输数据#

大多数程序的下一步是将数据传输到设备上。在PyCuda中，您将主要从 numpy 主机上的阵列。(但实际上，满足Python缓冲区接口的所有东西都可以工作，甚至 bytes 。)让我们制作一个4x4随机数组：：

import numpy
a = numpy.random.randn(4,4)

但是等待-- a 由双精度数字组成，但大多数nvidia设备只支持单精度：

a = a.astype(numpy.float32)

最后，我们需要一个地方来传输数据，所以我们需要在设备上分配内存：

a_gpu = cuda.mem_alloc(a.nbytes)

最后一步，我们需要将数据传输到GPU：

cuda.memcpy_htod(a_gpu, a)

执行内核#

对于本教程，我们将坚持一些简单的方法：编写代码将每个条目加倍 a_gpu . 为此，我们编写相应的CUDAC代码，并将其输入 pycuda.compiler.SourceModule ：：

mod = SourceModule("""
  __global__ void doublify(float *a)
  {
    int idx = threadIdx.x + threadIdx.y*4;
    a[idx] *= 2;
  }
  """)

如果没有任何错误，代码现在被编译并加载到设备上。我们找到了 pycuda.driver.Function 调用它，指定 a_gpu 作为参数，块大小为4x4:：

func = mod.get_function("doublify")
func(a_gpu, block=(4,4,1))

最后，我们从gpu取回数据并将其与原始数据一起显示 a ：：

a_doubled = numpy.empty_like(a)
cuda.memcpy_dtoh(a_doubled, a_gpu)
print(a_doubled)
print(a)

这将打印如下内容：

[[ 0.51360393  1.40589952  2.25009012  3.02563429]
 [-0.75841576 -1.18757617  2.72269917  3.12156057]
 [ 0.28826082 -2.92448163  1.21624792  2.86353827]
 [ 1.57651746  0.63500965  2.21570683 -0.44537592]]
[[ 0.25680196  0.70294976  1.12504506  1.51281714]
 [-0.37920788 -0.59378809  1.36134958  1.56078029]
 [ 0.14413041 -1.46224082  0.60812396  1.43176913]
 [ 0.78825873  0.31750482  1.10785341 -0.22268796]]

它奏效了！这就完成了我们的演练。谢天谢地，Pycuda接管了这里，为你做所有的清理工作，所以你完成了。不过，在下一节中，请继续阅读一些有价值的材料。

（你可以找到这个演示的代码 examples/demo.py 在pycuda源代码发行版中。）

显式内存副本的快捷方式#

这个 pycuda.driver.In ， pycuda.driver.Out 和 pycuda.driver.InOut 参数处理程序可以简化一些内存传输。例如，不是创建 a_gpu ，如果替换 a 很好，下面的代码可以使用：

func(cuda.InOut(a), block=(4, 4, 1))

准备好的调用#

使用内置的函数调用 pycuda.driver.Function.__call__() 方法导致类型标识的开销（请参见设备接口）为了获得与上面相同的效果而不产生这种开销，函数被绑定到参数类型（由python的标准库指定 struct 模块），然后调用。这也避免了使用 numpy.number 班级：

grid = (1, 1)
block = (4, 4, 1)
func.prepare("P")
func.prepared_call(grid, block, a_gpu)

好处：消除了并发症#

使用A pycuda.gpuarray.GPUArray 同样的效果可以用更少的写作来实现：

import pycuda.gpuarray as gpuarray
import pycuda.driver as cuda
import pycuda.autoinit
import numpy

a_gpu = gpuarray.to_gpu(numpy.random.randn(4,4).astype(numpy.float32))
a_doubled = (2*a_gpu).get()
print(a_doubled)
print(a_gpu)

高级主题#

结构#

（由董建华撰稿，代码见 examples/demo_struct.py ）

假设我们有以下结构，用于将多个可变长度数组加倍：

mod = SourceModule("""
    struct DoubleOperation {
        int datalen, __padding; // so 64-bit ptrs can be aligned
        float *ptr;
    };

    __global__ void double_array(DoubleOperation *a) {
        a = &a[blockIdx.x];
        for (int idx = threadIdx.x; idx < a->datalen; idx += blockDim.x) {
            a->ptr[idx] *= 2;
        }
    }
    """)

网格中的每个块（参见CUDA文档）将使其中一个数组加倍。这个 for 循环允许更多的数据元素比线程加倍，虽然效率不高，如果可以保证有足够数量的线程。接下来，创建结构的包装类，并实例化两个数组：

class DoubleOpStruct:
    mem_size = 8 + numpy.intp(0).nbytes
    def __init__(self, array, struct_arr_ptr):
        self.data = cuda.to_device(array)
        self.shape, self.dtype = array.shape, array.dtype
        packed_args = struct.pack("ixP", array.size, numpy.uintp(self.data))
        cuda.memcpy_htod(struct_arr_ptr, packed_args)

    def __str__(self):
        return str(cuda.from_device(self.data, self.shape, self.dtype))

struct_arr = cuda.mem_alloc(2 * DoubleOpStruct.mem_size)
do2_ptr = int(struct_arr) + DoubleOpStruct.mem_size

array1 = DoubleOpStruct(numpy.array([1, 2, 3], dtype=numpy.float32), struct_arr)
array2 = DoubleOpStruct(numpy.array([0, 4], dtype=numpy.float32), do2_ptr)
print("original arrays", array1, array2)

此代码使用 pycuda.driver.to_device() 和 pycuda.driver.from_device() 分配和复制值的功能，并演示如何使用分配内存块的偏移量。最后，代码可以被执行；下面演示了将两个数组加倍，然后只有第二个：

func = mod.get_function("double_array")
func(struct_arr, block = (32, 1, 1), grid=(2, 1))
print("doubled arrays", array1, array2)

func(numpy.intp(do2_ptr), block = (32, 1, 1), grid=(1, 1))
print("doubled second only", array1, array2, "\n")

使用CUDA阵列接口实现与其他库的互操作性#

内核调用可以传递来自其他CUDA库的数组 CUDA Array Interface 。例如，要将一个 CuPy 阵列：：

import cupy as cp

cupy_a = cp.random.randn(4, 4).astype(cp.float32)
func = mod.get_function("double_array")
func(cupy_a, block=(4, 4, 1), grid=(1, 1))

GPUArray 实现CUDA数组接口，因此它的实例可以从支持它的其他库传递到函数中。例如，若要将PyCUDA GPU阵列加倍，请使用 Numba 内核：：

from numba import cuda

a_gpu = gpuarray.to_gpu(numpy.random.randn(4, 4).astype(numpy.float32))

@cuda.jit
def double(x):
    i, j = cuda.grid(2)
    x[i, j] *= 2

double[(4, 4), (1, 1)](a_gpu)

从这里到哪里#

一旦你对基本知识有了足够的了解，就可以自由地深入了解设备接口 . 有关更多示例，请在 examples/ 分发的子目录。此文件夹还包含几个基准测试，以查看GPU和基于CPU的计算之间的差异。作为如何完成任务的参考，pycuda的测试套件 test/ 发行版的子目录可能也有帮助。