import pycuda.autoinit
import pycuda.driver as drv
import numpy
from pycuda.compiler import SourceModule
mod = SourceModule("""
#include <stdio.h>
__global__ void add_num(float *d_result, float *d_a, float *d_b)
{
const int i = threadIdx.x;
d_result[i] = d_a[i] + d_b[i];
}
""")
add_num = mod.get_function("add_num")
h_a = numpy.random.randn(1).astype(numpy.float32)
h_b = numpy.random.randn(1).astype(numpy.float32)
h_result = numpy.zeros_like(h_a)
d_a = drv.mem_alloc(h_a.nbytes)
d_b = drv.mem_alloc(h_b.nbytes)
d_result = drv.mem_alloc(h_result.nbytes)
drv.memcpy_htod(d_a,h_a)
drv.memcpy_htod(d_b,h_b)
add_num(
d_result, d_a, d_b,
block=(1,1,1), grid=(1,1))
drv.memcpy_dtoh(h_result,d_result)
print("Addition on GPU:")
print(h_a[0],"+", h_b[0] , "=" , h_result[0])
可使用PyCUDA中 driver类的mem_alloc函数来分配设备上的内存,用h_a.nbytes p函数找到内存的大小,并作为参数传递给函数。PyCUDA在driver类中提供了一个memcpy函数用于将数据从主机内存复制到设备显存,反之亦然。
用drv.memcpy_htod函数将数据从主机内存复制到设备显存,设备显存的指针作为第一个参数传递,主机内存的指针作为第二个参数传递。通过将设备指针、指定要启动的线程块、线程数的数字一起作为传递参数来调用add_num内核。在前面给出的代码中,使用一个线程启动一个块。最后通过drv.memcpy_dtoh函数将内核计算的执行结果数据复制回主机,然后显示在控制台上。
PyCUDA为内核调用提供一个更简单的API,不需要内存分配和内存复制,这是由API隐式完成的,可以通过使用PyCUDA中 driver类的In和Out函数来实现。修改后的数组加法代码如下:
import pycuda.autoinit
import pycuda.driver as drv
import numpy
N = 1
from pycuda.compiler import SourceModule
mod = SourceModule("""
__global__ void add_num(float *d_result, float *d_a, float *d_b)
{
const int i = threadIdx.x;
d_result[i] = d_a[i] + d_b[i];
}
""")
add_num = mod.get_function("add_num")
h_a = numpy.random.randn(N).astype(numpy.float32)
h_b = numpy.random.randn(N).astype(numpy.float32)
h_result = numpy.zeros_like(h_a)
add_num(
drv.Out(h_result), drv.In(h_a), drv.In(h_b),
block=(N,1,1), grid=(1,1))
print("Addition on GPU:")
for i in range(0,N):
print(h_a[i],"+", h_b[i] , "=" , h_result[i])
前面的代码中数组添加10个元素,而不是单个元素。内核函数与前面看到的代码完全相同,在主机上创建两个由10个随机数组成的数组,这次不再创建内存并将其传输到设备,而是直接调用内核,通过使用drv.out或drv.In指定数据的方向来修改内核函数,这简化了PyCUDA代码。
调用内核时,启动一个具有N个线程的线程块,这N个线程并行添加数组的N个元素,这将加速加法操作。内核函数计算结果通过drv.out 指令自动下载到主机内存中,因此该结果使用for循环直接打印在控制台上。