专栏首页Dechin的专栏Python实现GPU加速的基本操作

Python实现GPU加速的基本操作

技术背景

之前写过一篇讲述如何使用pycuda来在Python上写CUDA程序的博客。这个方案的特点在于完全遵循了CUDA程序的写法,只是支持了一些常用函数的接口,如果你需要自己写CUDA算子,那么就只能使用非常不Pythonic的写法。还有一种常见的方法是用cupy来替代numpy,相当于一个GPU版本的numpy。那么本文要讲述的是用numba自带的装饰器,来写一个非常Pythonic的CUDA程序。

CUDA的线程与块

GPU从计算逻辑来讲,可以认为是一个高并行度的计算阵列,我们可以想象成一个二维的像围棋棋盘一样的网格,每一个格子都可以执行一个单独的任务,并且所有的格子可以同时执行计算任务,这就是GPU加速的来源。那么刚才所提到的棋盘,每一列都认为是一个线程,并有自己的线程编号;每一行都是一个块,有自己的块编号。我们可以通过一些简单的程序来理解这其中的逻辑:

用GPU打印线程编号

# numba_cuda_test.py

from numba import cuda

@cuda.jit
def gpu():
    print ('threadIdx:', cuda.threadIdx.x)

if __name__ == '__main__':
    gpu[2,4]()
threadIdx: 0
threadIdx: 1
threadIdx: 2
threadIdx: 3
threadIdx: 0
threadIdx: 1
threadIdx: 2
threadIdx: 3

用GPU打印块编号

# numba_cuda_test.py

from numba import cuda

@cuda.jit
def gpu():
    print ('blockIdx:', cuda.blockIdx.x)

if __name__ == '__main__':
    gpu[2,4]()
blockIdx: 0
blockIdx: 0
blockIdx: 0
blockIdx: 0
blockIdx: 1
blockIdx: 1
blockIdx: 1
blockIdx: 1

用GPU打印块的维度

# numba_cuda_test.py

from numba import cuda

@cuda.jit
def gpu():
    print ('blockDim:', cuda.blockDim.x)

if __name__ == '__main__':
    gpu[2,4]()
blockDim: 4
blockDim: 4
blockDim: 4
blockDim: 4
blockDim: 4
blockDim: 4
blockDim: 4
blockDim: 4

用GPU打印线程的维度

# numba_cuda_test.py

from numba import cuda

@cuda.jit
def gpu():
    print ('gridDim:', cuda.gridDim.x)

if __name__ == '__main__':
    gpu[2,4]()
gridDim: 2
gridDim: 2
gridDim: 2
gridDim: 2
gridDim: 2
gridDim: 2
gridDim: 2
gridDim: 2

总结

我们可以用如下的一张图来总结刚才提到的GPU网格的概念,在上面的测试案例中,我们在GPU上划分一块2*4大小的阵列用于我们自己的计算,每一行都是一个块,每一列都是一个线程,所有的网格是同时执行计算的内容的(如果没有逻辑上的依赖的话)。

GPU所支持的最大并行度

我们可以用几个简单的程序来测试一下GPU的并行度,因为每一个GPU上的网格都可以独立的执行一个任务,因此我们认为可以分配多少个网格,就有多大的并行度。本机的最大并行应该是在

2^40

,因此假设我们给GPU分配

2^50

大小的网格,程序就会报错:

# numba_cuda_test.py

from numba import cuda

@cuda.jit
def gpu():
    pass

if __name__ == '__main__':
    gpu[2**50,1]()
    print ('Running Success!')

运行结果如下:

Traceback (most recent call last):
  File "numba_cuda_test.py", line 10, in <module>
    gpu[2**50,1]()
  File "/home/dechin/.local/lib/python3.7/site-packages/numba/cuda/compiler.py", line 822, in __call__
    self.stream, self.sharedmem)
  File "/home/dechin/.local/lib/python3.7/site-packages/numba/cuda/compiler.py", line 966, in call
    kernel.launch(args, griddim, blockdim, stream, sharedmem)
  File "/home/dechin/.local/lib/python3.7/site-packages/numba/cuda/compiler.py", line 699, in launch
    cooperative=self.cooperative)
  File "/home/dechin/.local/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py", line 2100, in launch_kernel
    None)
  File "/home/dechin/.local/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py", line 300, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "/home/dechin/.local/lib/python3.7/site-packages/numba/cuda/cudadrv/driver.py", line 335, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [1] Call to cuLaunchKernel results in CUDA_ERROR_INVALID_VALUE

而如果我们分配一个额定大小之内的网格,程序就可以正常的运行:

# numba_cuda_test.py

from numba import cuda

@cuda.jit
def gpu():
    pass

if __name__ == '__main__':
    gpu[2**30,1]()
    print ('Running Success!')

这里加了一个打印输出:

Running Success!

需要注意的是,两个维度上的可分配大小是不一致的,比如本机的上限是分配230*210大小的空间用于计算:

# numba_cuda_test.py

from numba import cuda

@cuda.jit
def gpu():
    pass

if __name__ == '__main__':
    gpu[2**30,2**10]()
    print ('Running Success!')

同样的,只要在允许的范围内都是可以执行成功的:

Running Success!

如果在本机上有多块GPU的话,还可以通过select_device的指令来选择执行指令的GPU编号:

# numba_cuda_test.py

from numba import cuda
cuda.select_device(1)
import time

@cuda.jit
def gpu():
    pass

if __name__ == '__main__':
    gpu[2**30,2**10]()
    print ('Running Success!')

如果两块GPU的可分配空间一致的话,就可以运行成功:

Running Success!

GPU的加速效果

前面我们经常提到一个词叫GPU加速,GPU之所以能够实现加速的效果,正源自于GPU本身的高度并行性。这里我们直接用一个数组求和的案例来说明GPU的加速效果,这个案例需要得到的结果是

b_j=a_j+b_j

,将求和后的值赋值在其中的一个输入数组之上,以节省一些内存空间。当然,如果这个数组还有其他的用途的话,是不能这样操作的。具体代码如下:

# gpu_add.py

from numba import cuda
cuda.select_device(1)
import numpy as np
import time

@cuda.jit
def gpu(a,b,DATA_LENGHTH):
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
    if idx < DATA_LENGHTH:
        b[idx] += a[idx]

if __name__ == '__main__':
    np.random.seed(1)
    DATA_EXP_LENGTH = 20
    DATA_DIMENSION = 2**DATA_EXP_LENGTH
    np_time = 0.0
    nb_time = 0.0
    for i in range(100):
        a = np.random.randn(DATA_DIMENSION).astype(np.float32)
        b = np.random.randn(DATA_DIMENSION).astype(np.float32)
        a_cuda = cuda.to_device(a)
        b_cuda = cuda.to_device(b)
        time0 = time.time()
        gpu[DATA_DIMENSION,4](a_cuda,b_cuda,DATA_DIMENSION)
        time1 = time.time()
        c = b_cuda.copy_to_host()
        time2 = time.time()
        d = np.add(a,b)
        time3 = time.time()
        if i == 0:
            print ('The error between numba and numpy is: ', sum(c-d))
            continue
        np_time += time3 - time2
        nb_time += time1 - time0
    print ('The time cost of numba is: {}s'.format(nb_time))
    print ('The time cost of numpy is: {}s'.format(np_time))

需要注意的是,基于Numba实现的Python的GPU加速程序,采用的jit即时编译的模式,也就是说,在运行调用到相关函数时,才会对其进行编译优化。换句话说,第一次执行这一条指令的时候,事实上达不到加速的效果,因为这个运行的时间包含了较长的一段编译时间。但是从第二次运行调用开始,就不需要重新编译,这时候GPU加速的效果就体现出来了,运行结果如下:

$ python3 gpu_add.py The error between numba and numpy is:  0.0
The time cost of numba is: 0.018711328506469727s
The time cost of numpy is: 0.09502553939819336s

可以看到,即使是相比于Python中优化程度十分强大的的Numpy实现,我们自己写的GPU加速的程序也能够达到5倍的加速效果(在前面一篇博客中,针对于特殊计算场景,加速效果可达1000倍以上),而且可定制化程度非常之高。

总结概要

本文针对于Python中使用Numba的GPU加速程序的一些基本概念和实现的方法,比如GPU中的线程和模块的概念,以及给出了一个矢量加法的代码案例,进一步说明了GPU加速的效果。需要注意的是,由于Python中的Numba实现是一种即时编译的技术,因此第一次运算时的时间会明显较长,所以我们一般说GPU加速是指从第二步开始的运行时间。对于一些工业和学界常见的场景,比如分子动力学模拟中的系统演化,或者是深度学习与量子计算中的参数优化,都是相同维度参数多步运算的一个过程,非常适合使用即时编译的技术,配合以GPU高度并行化的加速效果,能够在实际工业和学术界的各种场景下发挥巨大的作用。

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

我来说两句

0 条评论
登录 后参与评论

相关文章

  • Python的基本操作

    (1)单引号、双引号 (2)三引号 Python没有多行注释,使用#进行单行注释,可以使用三引号模拟实现多行注释。 (3)转义符

    魏晓蕾
  • Nginx零成本、易操作实现网站视频加速

    最近有朋友做了个网站,网站上面放了几个视频,觉得视频播放慢、卡顿,想让我帮忙优化下,今天就介绍个零成本,操作又简单的方法优化网站上视频播放速度

    李俊鹏
  • 小蛇学python(22)pytorch配置cuda实现GPU加速

    深度学习如火如荼,使用普通的cpu来跑模型真的让人急死,就算最普通的垃圾显卡,只要支持cuda,就可以实现gpu加速,其速度至少是cpu的5倍。

    用户2145057
  • Java/C/Python_实现顺序表的基本操作

    详细源码参考地址: https://github.com/jackaroo2020/my-algorithm

    Java架构师必看
  • Python3实现打格点算法的GPU加速

    在数学和物理学领域,总是充满了各种连续的函数模型。而当我们用现代计算机的技术去处理这些问题的时候,事实上是无法直接处理连续模型的,绝大多数的情况下都要转化成一个...

    DechinPhy
  • TensorFlow使用Graph的基本操作的实现

    在tensorflow中,一个程序默认是建立一个图的,除了系统自动建立图以外,我们还可以手动建立图,并做一些其他的操作。

    砸漏
  • python对kafka的基本操作

    from kafka import KafkaProducer from kafka import KafkaConsumer from kafka.struc...

    py3study
  • NVIDIA的python-GPU算法生态 ︱ RAPIDS 0.10

    随着新版本的推出,RAPIDS 迎来了其推出一周年纪念日。回顾所经历的一年,RAPIDS团队就社区对该项目的关心和支持表示衷心的感谢。此前,RAPIDS获得了其...

    素质
  • 栈与栈的实现栈栈的基本操作栈的实现

    栈 栈是一种基础的数据结构,只从一端读写数据。基本特点就”后进先出“,例如顺序入栈1,2,3,4,5,再顺序出栈是5,4,3,2,1 栈的基本操作 栈的基本操作...

    月见樽
  • Python|处理word的基本操作

    众所周知python有很多第三方库,这也是python简单实用的原因。要想用python处理word文档就需要安装python-docx库。

    算法与编程之美
  • 【python-opencv】图像的基本操作

    你可以通过行和列坐标来访问像素值。对于 BGR 图像,它返回一个由蓝色、绿色和红色值组成的数组。对于灰度图像,只返回相应的灰度。

    西西嘛呦
  • 如何在 GPU 上加速数据科学

    数据科学家需要算力。无论您是用 pandas 处理一个大数据集,还是用 Numpy 在一个大矩阵上运行一些计算,您都需要一台强大的机器,以便在合理的时间内完成这...

    AI研习社
  • 使用python脚本实现mysql误操作

    1.简介 在oracle数据库中,当一个误操作被提交后,我们可以通过oracle提供的闪回功能将表闪回至误操作之前的状态。mysql中没有原生的flushbac...

    py3study
  • 如何在 GPU 上加速数据科学

    我们认为使用大型模型架构和相同数据在XLNet 和BERT之间进行公平的比较研究具有重要的科学价值。

    AI科技评论
  • 免费!Google Colab现已支持英伟达T4 GPU

    【新智元导读】Google Colab现在提供免费的T4 GPU。Colab是Google的一项免费云端机器学习服务,T4GPU耗能仅为70瓦,是面向现有数...

    新智元
  • 深度学习之在 Ubuntu 上安装 Keras 及其依赖

    Keras是一个由Python编写的开源人工神经网络库,可以作为Tensorflow、Microsoft-CNTK和Theano的高阶应用程序接口,进行深度学习...

    李小白是一只喵
  • 很火的深度学习框架PyTorch怎么用?手把手带你安装配置

    PyTorch是Facebook团队于2017年1月发布的一个深度学习框架,虽然晚于TensorFlow、Keras等框架,但自发布之日起,其关注度就在不断上升...

    CDA数据分析师
  • Reddit热议:为什么PyTorch比TensorFlow更快?

    近日,Reddit 上有一个热帖:为什么 PyTorch 和 TensorFlow 一样快 (有时甚至比 TensorFlow 更快)?

    新智元
  • Reddit热议:为什么PyTorch比TensorFlow更快?

    近日,Reddit 上有一个热帖:为什么 PyTorch 和 TensorFlow 一样快 (有时甚至比 TensorFlow 更快)?

    代码医生工作室

扫码关注云+社区

领取腾讯云代金券