pycuda之流类

持续创作，加速成长！这是我参与「掘金日新计划 · 10 月更文挑战」的第30天，点击查看活动详情

先解释一下，什么是cuda的流，其实CUDA的流就是在GPU上按照顺序运行的一系列操作。但是，仅使用单个CUDA的流是无法获取我们所需要的并发性，因此如果想要获得更高的并发性，必须让主机向GPU发出多个CUDA流，这句话的意思就是说我们应该主动让CPU交替的启动不同的CUDA流的GPU操作，以便能够充分的利用当前的GPU流。

下面，本博客通过使用GPU流和不使用GPU流两个方面对程序的运行速度进行分析，看一下使用了GPU流之后速度是否加快，以及加快的速度大概是原来的多少倍。本例通过对随机生成的一个向量进行乘法和除法操作，而且乘法和除法分别为互逆的操作，最后通过numpy库的allclose函数进行检验。

首先，看一下不使用GPU流时的代码

import pycuda
import pycuda.autoinit
import pycuda.driver  as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
import numpy as np
import time


cuda_kernal = SourceModule(
    """
    __global__ void test_stream(float *array, int array_len)
    {
        int cu_id = threadIdx.x + blockDim.x * blockIdx.x;
        int iter_num = array_len / blockDim.x;
        for(int j=0; j<iter_num; j++)
        {
            int i = j * blockDim.x + cu_id;
            for(int k=0; k<50; k++)
            {
                array[i] *= 2.0;
                array[i] /= 2.0;
            }
        } 
    }
    """
)

multi_kernal = cuda_kernal.get_function("test_stream")
in_data = []
gpu_data = []
out_data = []
num_array = 200
array_len = 1024**2

for _ in range(num_array):
    # 注意这里，使用流时需要添加一条
    in_data.append(np.random.randn(array_len).astype("float32"))

t1 = time.time()
for i in range(num_array):
    # # 注意这里，使用流时需要替换
    gpu_data.append(gpuarray.to_gpu(in_data[i]))

for i in range(num_array):
    # # 注意这里，使用流时需要替换
    multi_kernal(gpu_data[i], np.int32(array_len), block=(64, 1, 1), grid=(1, 1, 1))

for i in range(num_array):
    # # 注意这里，使用流时需要替换
    out_data.append(gpu_data[i].get())

t2 = time.time()

for i in range(num_array):
    assert (np.allclose(out_data[i], in_data[i]))

print("time used:", t2-t1)

注意，代码中的1024**2，不要写的太大，不然可能会出现下面的错误，当出现下面的错误时不要着急，将此处的数值改小即可，经过测试，我将数值置3时就报错了。

最后正常情况下执行的结果如下，程序执行的打印结果与GPU的性能相关。可以看到，在本实验环境下(RTX2070，笔记本)，执行上述代码需要2.1089134216308594秒。鉴于使用PyCUDA流的代码改动不大，因此对上述代码进行稍微的改动，改动后的代码如下。

按照上面代码的注释部分，对下面的代码进行替换即可。

stream_data = []
for _ in range(num_array):
    in_data.append(np.random.randn(array_len).astype("float32"))
    stream_data.append(drv.Stream())

for k in range(num_array):
    gpu_data.append(gpuarray.to_gpu_async(in_data[k], stream=stream_data[k]))

for i in range(num_array):
    multi_kernal(gpu_data[i], np.int32(array_len), block=(64, 1, 1), grid=(1, 1, 1), stream=stream_data[i])

for i in range(num_array):
    out_data.append(gpu_data[i].get_async(stream=stream_data[i]))

执行后的结果如下

可以看到改进后的代码相较于原来的代码快了4.655倍。

经过上述的实验结果比较，同样都使用GPU并行处理的情况下，使用流可以相较于原来快4.6倍，这样的处理速度对于处理大批量的数据时且在核函数变化不大的情况下效果较好。