目錄
- CUDA的線程與塊
- 用GPU打印線程編號(hào)
- 用GPU打印塊編號(hào)
- 用GPU打印塊的維度
- 用GPU打印線程的維度
- 總結(jié)
- GPU所支持的最大并行度
- GPU的加速效果
- 總結(jié)概要
CUDA的線程與塊
GPU從計(jì)算邏輯來講,可以認(rèn)為是一個(gè)高并行度的計(jì)算陣列,我們可以想象成一個(gè)二維的像圍棋棋盤一樣的網(wǎng)格,每一個(gè)格子都可以執(zhí)行一個(gè)單獨(dú)的任務(wù),并且所有的格子可以同時(shí)執(zhí)行計(jì)算任務(wù),這就是GPU加速的來源。那么剛才所提到的棋盤,每一列都認(rèn)為是一個(gè)線程,并有自己的線程編號(hào);每一行都是一個(gè)塊,有自己的塊編號(hào)。我們可以通過一些簡(jiǎn)單的程序來理解這其中的邏輯:
用GPU打印線程編號(hào)
# 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打印塊編號(hào)
# 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
總結(jié)
我們可以用如下的一張圖來總結(jié)剛才提到的GPU網(wǎng)格的概念,在上面的測(cè)試案例中,我們?cè)贕PU上劃分一塊2*4大小的陣列用于我們自己的計(jì)算,每一行都是一個(gè)塊,每一列都是一個(gè)線程,所有的網(wǎng)格是同時(shí)執(zhí)行計(jì)算的內(nèi)容的(如果沒有邏輯上的依賴的話)。
GPU所支持的最大并行度
我們可以用幾個(gè)簡(jiǎn)單的程序來測(cè)試一下GPU的并行度,因?yàn)槊恳粋€(gè)GPU上的網(wǎng)格都可以獨(dú)立的執(zhí)行一個(gè)任務(wù),因此我們認(rèn)為可以分配多少個(gè)網(wǎng)格,就有多大的并行度。本機(jī)的最大并行應(yīng)該是在\(2^40\),因此假設(shè)我們給GPU分配\(2^50\)大小的網(wǎng)格,程序就會(huì)報(bào)錯(cuò):
# numba_cuda_test.py
from numba import cuda
@cuda.jit
def gpu():
pass
if __name__ == '__main__':
gpu[2**50,1]()
print ('Running Success!')
運(yùn)行結(jié)果如下:
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
而如果我們分配一個(gè)額定大小之內(nèi)的網(wǎng)格,程序就可以正常的運(yùn)行:
# numba_cuda_test.py
from numba import cuda
@cuda.jit
def gpu():
pass
if __name__ == '__main__':
gpu[2**30,1]()
print ('Running Success!')
這里加了一個(gè)打印輸出:
Running Success!
需要注意的是,兩個(gè)維度上的可分配大小是不一致的,比如本機(jī)的上限是分配230*210大小的空間用于計(jì)算:
# numba_cuda_test.py
from numba import cuda
@cuda.jit
def gpu():
pass
if __name__ == '__main__':
gpu[2**30,2**10]()
print ('Running Success!')
同樣的,只要在允許的范圍內(nèi)都是可以執(zhí)行成功的:
Running Success!
如果在本機(jī)上有多塊GPU的話,還可以通過select_device
的指令來選擇執(zhí)行指令的GPU編號(hào):
# 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的可分配空間一致的話,就可以運(yùn)行成功:
Running Success!
GPU的加速效果
前面我們經(jīng)常提到一個(gè)詞叫GPU加速,GPU之所以能夠?qū)崿F(xiàn)加速的效果,正源自于GPU本身的高度并行性。這里我們直接用一個(gè)數(shù)組求和的案例來說明GPU的加速效果,這個(gè)案例需要得到的結(jié)果是\(b_j=a_j+b_j\),將求和后的值賦值在其中的一個(gè)輸入數(shù)組之上,以節(jié)省一些內(nèi)存空間。當(dāng)然,如果這個(gè)數(shù)組還有其他的用途的話,是不能這樣操作的。具體代碼如下:
# 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實(shí)現(xiàn)的Python的GPU加速程序,采用的jit即時(shí)編譯的模式,也就是說,在運(yùn)行調(diào)用到相關(guān)函數(shù)時(shí),才會(huì)對(duì)其進(jìn)行編譯優(yōu)化。換句話說,第一次執(zhí)行這一條指令的時(shí)候,事實(shí)上達(dá)不到加速的效果,因?yàn)檫@個(gè)運(yùn)行的時(shí)間包含了較長(zhǎng)的一段編譯時(shí)間。但是從第二次運(yùn)行調(diào)用開始,就不需要重新編譯,這時(shí)候GPU加速的效果就體現(xiàn)出來了,運(yùn)行結(jié)果如下:
$ 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中優(yōu)化程度十分強(qiáng)大的的Numpy實(shí)現(xiàn),我們自己寫的GPU加速的程序也能夠達(dá)到5倍的加速效果(在前面一篇博客中,針對(duì)于特殊計(jì)算場(chǎng)景,加速效果可達(dá)1000倍以上),而且可定制化程度非常之高。
總結(jié)概要
本文針對(duì)于Python中使用Numba的GPU加速程序的一些基本概念和實(shí)現(xiàn)的方法,比如GPU中的線程和模塊的概念,以及給出了一個(gè)矢量加法的代碼案例,進(jìn)一步說明了GPU加速的效果。需要注意的是,由于Python中的Numba實(shí)現(xiàn)是一種即時(shí)編譯的技術(shù),因此第一次運(yùn)算時(shí)的時(shí)間會(huì)明顯較長(zhǎng),所以我們一般說GPU加速是指從第二步開始的運(yùn)行時(shí)間。對(duì)于一些工業(yè)和學(xué)界常見的場(chǎng)景,比如分子動(dòng)力學(xué)模擬中的系統(tǒng)演化,或者是深度學(xué)習(xí)與量子計(jì)算中的參數(shù)優(yōu)化,都是相同維度參數(shù)多步運(yùn)算的一個(gè)過程,非常適合使用即時(shí)編譯的技術(shù),配合以GPU高度并行化的加速效果,能夠在實(shí)際工業(yè)和學(xué)術(shù)界的各種場(chǎng)景下發(fā)揮巨大的作用。
到此這篇關(guān)于Python實(shí)現(xiàn)GPU加速的基本操作的文章就介紹到這了,更多相關(guān)Python GPU加速內(nèi)容請(qǐng)搜索腳本之家以前的文章或繼續(xù)瀏覽下面的相關(guān)文章希望大家以后多多支持腳本之家!
您可能感興趣的文章:- Python基于pyCUDA實(shí)現(xiàn)GPU加速并行計(jì)算功能入門教程