GPU编程（基于Python和CUDA）（三）——逐元素运算核（ElementwiseKernel）

系列文章目录

GPU编程（基于Python和CUDA）（一）——零基础安装pycuda
GPU编程（基于Python和CUDA）（二）——显示GPU信息
GPU编程（基于Python和CUDA）（三）——逐元素运算核（ElementwiseKernel）
GPU编程（基于Python和CUDA）（四）——Mandelbort集

---

---

前言

在向量运算和矩阵运算中，对于向量和矩阵按照元素逐个运算十分常见，在本篇中将采用向量与标量相乘为例子介绍逐元素运算核

标量与向量相乘

常用的标量与向量相乘的方法有直接相乘核使用逐元素运算核相乘，在介绍相乘之前我们还需要了解一些准备工作

预备知识

• pycuda初始化：pycuda是需要先初始化才能正常使用的，一般来讲会使用import pycuda.autoinit进行初始化，需要注意的是在编辑器上这一行代码会显示未被使用，以pycharm为例，显示是灰色的
• 将数据传到GPU：我们使用from pycuda import gpuarray引入gpuarray，他的作用和numpy中的array相似，但他的运行是依赖GPU的，通过gpuarray.to_gpu可将数据传向GPU，通过get从GPU取回数据

直接相乘

在对比逐元素运算核的时候我们自然不能仅仅与列表逐元素相乘做对比，那样逐元素运算核的优势太明显了，而应该与使用CPU的numpy作比较。这里对比的逻辑是如果还不如使用numpy速度快，那还不如不进行GPU编程，直接使用numpy就可以了。

1.导入包
如果没有安装numpy包可用使用pip install numpy进行安装

import numpy as np
import pycuda.autoinit
from pycuda import gpuarray
from time import time

2.测试numpy时间
我们使用numpy生成一个随机向量host_data，使用time来获取时间

host_data = np.float32(np.random.random(5000000))

t1 = time()
host_data_x2 = host_data * np.float32(2)
t2 = time()

print("CPU总耗时:%f" % (t2-t1))

总耗时的结果（不同设备的运行结果可能不同）如下：

CPU总耗时:0.003995

3.测试gpuarray时间
我们将host_data传入GPU后进行运算，然后在将数据从GPU取出，测算时间后使用np.allclose对比两种方式的运行结果是否一致

device_data = gpuarray.to_gpu(host_data)

t1 = time()
device_data_x2 = device_data * np.float32(2)
t2 = time()

from_device = device_data_x2.get()
print("GPU总耗时:%f" % (t2-t1))

print("结果是否一致?:{}".format(np.allclose(from_device, host_data_x2)))

结果（不同设备的运行结果可能不同）如下：

GPU总耗时:1.321204
结果是否一致?:True

我们看到，时间的运行结果是GPU的耗时会比CPU的耗时还要高，这是为什么呢？因为在pycuda首次使用GPU内核函数时需要使用nvcc编译器编译CUDA C代码，也就是说，在首次调用逐元素运算核时会进行编译，这段时间耗时比较就，那么我们可以得出结论，如果多次调用就会更快了，我们下边来验证我们的猜想。
4.验证
首先我们把上述3个步骤的代码按照顺序放到py文件中，这里笔者的py文件的名称为time_calc.py
随后为了更方便的验证，我们可以在jupyter notebook中使用run time_calc.py调用这个py文件。多次调用观察结果
如果没有jupyter notebook，也可以把相乘和计算时间的代码多复制几遍观察结果。
两次计算结果（不同设备的运行结果可能不同)如下：

CPU总耗时:0.004980
GPU总耗时:1.182359
结果是否一致?:True
CPU总耗时:0.003993
GPU总耗时:0.000000
结果是否一致?:True

使用逐元素运算核相乘

1.导入包
和之前相比仅仅多了一个ElementwiseKernel的导入

import numpy as np
import pycuda.autoinit
from pycuda import gpuarray
from time import time
from pycuda.elementwise import ElementwiseKernel

2.定义CUDA C内核函数
整个定义过程和C/C++的函数定义类似，只不过ElementwiseKernel增加了自动遍历的功能。

• 首先第一行float *in, float *out是输入和输出，因为我们输入和输出都是一个向量，对应到C/C++就是数组，所以这里传入指针类型
• 第二行是遍历过程，因为ElementwiseKernel可以自动遍历，所以不需要自己写循环语句，只需要按照遍历格式写运算过程即可
• 第三行是CUDA C内核函数名，一般可以随意设置，但建议设置的通俗易懂一些

gpu_2x_ker = ElementwiseKernel(
    "float *in, float *out",  # 设置输入变量和输出变量的格式（指针形式）
    "out[i] = 2 * in[i];",  # i为索引，pycuda会自动设置
    "gpu_2x_ker",  # CUDA C内核函数名
)

3.测试
这次测试和之前有所不同，除了将数据送如GPU外还应该使用gpuarray.empty_like新建一个空数组。随后就可以使用我们定义好的内核gpu_2x_ker了

host_data = np.float32(np.random.random(5000000))
t1 = time()
host_data_2x = host_data * np.float32(2)
t2 = time()
print("CPU耗时:%f" % (t2-t1))
device_data = gpuarray.to_gpu(host_data)
device_data_2x = gpuarray.empty_like(device_data)
t1 = time()
gpu_2x_ker(device_data, device_data_2x)
t2 = time()
from_device = device_data_2x.get()
print("GPU耗时:%f" % (t2-t1))
print("结果是否一致?:{}".format(np.allclose(from_device, host_data_2x)))

同样的，也是第一次使用的时候会很慢，但这里要注意的是，因为内核是自己定义的， 如果使用jupyter notebook调用会重复定义并编译，并不能达到提速的目的，这里只能将测试部分的代码复制一遍再运行，需要注意的是，只复制测试部分，定义内核函数的代码不要动。

测试结果（不同设备的运行结果可能不同)如下：

CPU耗时:0.004982
GPU耗时:1.423676
CPU耗时:0.005979
GPU耗时:0.000000
结果是否一致?:True