PyCUDA内核函数

元素级内核函数

import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
from pycuda.elementwise import ElementwiseKernel

add = ElementwiseKernel(
        "float *d_a, float *d_b, float *d_c",
        "d_c[i] = d_a[i] + d_b[i]",
        "add")

# create a couple of random matrices with a given shape
from pycuda.curandom import rand as curand
shape = 1000000
d_a = curand(shape)
d_b = curand(shape)
d_c = gpuarray.empty_like(d_a)
start = drv.Event()
end = drv.Event()
start.record()
add(d_a, d_b, d_c)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Addition of %d element of GPU"%shape)
print("%fs" % (secs))
# check the result
if d_c == (d_a + d_b):
    print("The sum computed on GPU is correct")

使用pycuda.elementwise.ElementwiseKernel函数去定义元素相关的内核函数，它需要3个参数:第一个参数是内核函数的参数列表;第二个参数定义对每个元素执行的操作;第三个参数指定内核函数的名称。
两个数组被pycuda.curandom类的curand 函数初始化为随机数，这还是一个有用的功能，可以消除在主机上初始化然后上传到设备显存的需要，然后创建一个空的GPU数组来存储结果。将这三个变量作为参数传递来调用add内核。使用CUDA事件计算添加100万个元素所需的时间，并显示在控制台上。

规约内核函数

归约运算可以定义为使用某些表达式将元素集合缩减为单个值，这在各种并行计算应用中非常有用。以向量间点乘的计算为例说明PyCUDA的归约概念。

import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
from pycuda.reduction import ReductionKernel
import pycuda.autoinit

n=5
start = drv.Event()
end = drv.Event()
start.record()
d_a = gpuarray.arange(n,dtype= numpy.uint32)
d_b = gpuarray.arange(n,dtype= numpy.uint32)
kernel = ReductionKernel(numpy.uint32,neutral="0",reduce_expr="a+b",map_expr="d_a[i]*d_b[i]",arguments="int *d_a,int *d_b")
d_result = kernel(d_a,d_b).get()
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Vector A")
print(d_a)
print("Vector B")
print(d_b)
print("The computed dot product using reduction:")
print(d_result)
print("Dot Product on GPU")
print("%fs" % (secs))

PyCUDA提供pycuda.reducation.ReductionKernel类来定义归约内核函数，它需要很多参数:第一个参数是输出的数据类型;第二个参数是中性的，通常定义为0;第三个参数是用于缩减元素集合的表达式，添加操作在前面的代码中定义;第四个参数定义为在归约前用于操作数之间映射操作的表达式，在代码中定义了元素乘法;最后一个参数定义了内核函数的参数。
计算点乘的归约内核函数需要两个向量之间的元素相乘，然后求和。使用arange函数定义两个向量，其工作方式与Python中的 range函数类似，但arange将数组保存在设备上。通过将这两个向量作为参数传递来调用内核函数，结果到主机获取。使用CUDA事件计算所需的时间，并与点乘结果一起显示在控制台上。

scan内核函数

scan 操作是另一个非常重要的并行计算范例，指派特定函数去计算输人序列的第一项,再将这个计算结果与输人序列的第二项作为输入提供给这个特定函数计算，所有的中间计算结果形成输出的序列。这个概念可以用于各种应用，下面以累加为例，说明PyCUDA中scan内核函数的概念。

import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
from pycuda.scan import InclusiveScanKernel
import pycuda.autoinit

n=10
start = drv.Event()
end = drv.Event()
start.record()
kernel = InclusiveScanKernel(numpy.uint32,"a+b")
h_a = numpy.random.randint(1,10,n).astype(numpy.int32)
d_a = gpuarray.to_gpu(h_a)
kernel(d_a)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
assert(d_a.get() == numpy.cumsum(h_a,axis=0)).all()
print("The input data:")
print(h_a)
print("The computed cumulative sum using Scan:")
print(d_a.get())
print("Cumulative Sum on GPU")
print("%fs" % (secs))

PyCUDA提供了pycuda.scan.InclusiveScanKernel类来定义包含scan 操作的内核函数,要求以输出与扫描操作的数据类型为参数，为这个累加和的函数指定加法运算，以随机整数的数组作为此内核函数的输人，内核输出将与输入具有相同的大小。输入和输出向量以及计算累加和，以及使用的时间显示在控制台上。