Python實現GPU加速的基本操作

Posted on by

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高度並行化的加速效果，能夠在實際工業和學術界的各種場景下發揮巨大的作用。

到此這篇關於Python實現GPU加速的基本操作的文章就介紹到這瞭,更多相關Python GPU加速內容請搜索WalkonNet以前的文章或繼續瀏覽下面的相關文章希望大傢以後多多支持WalkonNet！

推薦閱讀：