教程

入门

在使用PycUDA之前,必须导入并初始化:

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

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

传输数据

大多数程序的下一步是将数据传输到设备上。在pycuda中,您将主要从 numpy 主机上的阵列。(但实际上,满足python缓冲区接口的所有东西都将工作,甚至 str )让我们制作一个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.Inpycuda.driver.Outpycuda.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
        cuda.memcpy_htod(int(struct_arr_ptr), numpy.getbuffer(numpy.int32(array.size)))
        cuda.memcpy_htod(int(struct_arr_ptr) + 8, numpy.getbuffer(numpy.intp(int(self.data))))
    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")

从这里到哪里

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