1. 初始化，并將必要的數(shù)據(jù)拷貝到GPU設(shè)備的顯存上。
2. CPU調(diào)用GPU函數(shù)，啟動(dòng)GPU多個(gè)核心同時(shí)進(jìn)行計(jì)算。
3. CPU與GPU異步計(jì)算。
4. 將GPU計(jì)算結(jié)果拷貝回主機(jī)端，得到計(jì)算結(jié)果。

一個(gè)名為gpu_print.py的GPU程序如下所示：

from numba import cuda

def cpu_print():
    print("print by cpu.")

@cuda.jit
def gpu_print():
    # GPU核函數(shù)
    print("print by gpu.")

def main():
    gpu_print[1, 2]()
    cuda.synchronize()
    cpu_print()

if __name__ == "__main__":
    main()

使用CUDA_VISIBLE_DEVICES=‘0’ python gpu_print.py執(zhí)行這段代碼，得到的結(jié)果為：

print by gpu.
print by gpu.
print by cpu.

與傳統(tǒng)的Python CPU代碼不同的是：

• 使用from numba import cuda引入cuda庫(kù)
• 在GPU函數(shù)上添加@cuda.jit裝飾符，表示該函數(shù)是一個(gè)在GPU設(shè)備上運(yùn)行的函數(shù)，GPU函數(shù)又被稱為核函數(shù)。
• 主函數(shù)調(diào)用GPU核函數(shù)時(shí)，需要添加如[1,2]這樣的執(zhí)行配置，這個(gè)配置是在告知GPU以多大的并行粒度同時(shí)進(jìn)行計(jì)算。gpu_print1,2表示同時(shí)開啟2個(gè)線程并行地執(zhí)行g(shù)pu_print函數(shù)，函數(shù)將被并行地執(zhí)行2次。下文會(huì)深入探討如何設(shè)置執(zhí)行配置。
• GPU核函數(shù)的啟動(dòng)方式是異步的：?jiǎn)?dòng)GPU函數(shù)后，CPU不會(huì)等待GPU函數(shù)執(zhí)行完畢才執(zhí)行下一行代碼。必要時(shí)，需要調(diào)用cuda.synchronize()，告知CPU等待GPU執(zhí)行完核函數(shù)后，再進(jìn)行CPU端后續(xù)計(jì)算。這個(gè)過(guò)程被稱為同步，也就是GPU執(zhí)行流程圖中的紅線部分。如果不調(diào)用cuda.synchronize()函數(shù)，執(zhí)行結(jié)果也將改變，"print bycpu.將先被打印。雖然GPU函數(shù)在前，但是程序并沒(méi)有等待GPU函數(shù)執(zhí)行完，而是繼續(xù)執(zhí)行后面的cpu_print函數(shù)，由于CPU調(diào)用GPU有一定的延遲，反而后面的cpu_print先被執(zhí)行，因此cpu_print的結(jié)果先被打印了出來(lái)。

Thread層次結(jié)構(gòu)

前面的程序中，核函數(shù)被GPU并行地執(zhí)行了2次。在進(jìn)行GPU并行編程時(shí)需要定義執(zhí)行配置來(lái)告知以怎樣的方式去并行計(jì)算，比如上面打印的例子中，是并行地執(zhí)行2次，還是8次，還是并行地執(zhí)行20萬(wàn)次，或者2000萬(wàn)次。2000萬(wàn)的數(shù)字太大，遠(yuǎn)遠(yuǎn)多于GPU的核心數(shù)，如何將2000萬(wàn)次計(jì)算合理分配到所有GPU核心上。解決這些問(wèn)題就需要弄明白CUDA的Thread層次結(jié)構(gòu)。

CUDA將核函數(shù)所定義的運(yùn)算稱為線程（Thread），多個(gè)線程組成一個(gè)塊（Block），多個(gè)塊組成網(wǎng)格（Grid）。這樣一個(gè)grid可以定義成千上萬(wàn)個(gè)線程，也就解決了并行執(zhí)行上萬(wàn)次操作的問(wèn)題。例如，把前面的程序改為并行執(zhí)行8次：可以用2個(gè)block，每個(gè)block中有4個(gè)thread。原來(lái)的代碼可以改為gpu_print2, 4，其中方括號(hào)中第一個(gè)數(shù)字表示整個(gè)grid有多少個(gè)block，方括號(hào)中第二個(gè)數(shù)字表示一個(gè)block有多少個(gè)thread。

實(shí)際上，線程（thread）是一個(gè)編程上的軟件概念。從硬件來(lái)看，thread運(yùn)行在一個(gè)CUDA核心上，多個(gè)thread組成的block運(yùn)行在Streaming Multiprocessor（SM的概念詳見本系列第一篇文章），多個(gè)block組成的grid運(yùn)行在一個(gè)GPU顯卡上。

CUDA提供了一系列內(nèi)置變量，以記錄thread和block的大小及索引下標(biāo)。以[2, 4]這樣的配置為例：blockDim.x變量表示block的大小是4，即每個(gè)block有4個(gè)thread，threadIdx.x變量是一個(gè)從0到blockDim.x - 1（4-1=3）的索引下標(biāo)，記錄這是第幾個(gè)thread；gridDim.x變量表示grid的大小是2，即每個(gè)grid有2個(gè)block，blockIdx.x變量是一個(gè)從0到gridDim.x - 1（2-1=1）的索引下標(biāo)，記錄這是第幾個(gè)block。

某個(gè)thread在整個(gè)grid中的位置編號(hào)為：threadIdx.x + blockIdx.x * blockDim.x。

利用上述變量，我們可以把之前的代碼豐富一下：

from numba import cuda

def cpu_print(N):
    for i in range(0, N):
        print(i)

@cuda.jit
def gpu_print(N):
    idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x 
    if (idx < N):
        print(idx)

def main():
    print("gpu print:")
    gpu_print[2, 4](8)
    cuda.synchronize()
    print("cpu print:")
    cpu_print(8)

if __name__ == "__main__":
    main()

執(zhí)行結(jié)果為：

gpu print:
0
1
2
3
4
5
6
7
cpu print:
0
1
2
3
4
5
6
7

這里的GPU函數(shù)在每個(gè)CUDA thread中打印了當(dāng)前thread的編號(hào)，起到了CPU函數(shù)for循環(huán)同樣的作用。因?yàn)閒or循環(huán)中的計(jì)算內(nèi)容互相不依賴，也就是說(shuō)，某次循環(huán)只是專心做自己的事情，循環(huán)第i次不影響循環(huán)第j次的計(jì)算，所以這樣互相不依賴的for循環(huán)非常適合放到CUDA thread里做并行計(jì)算。在實(shí)際使用中，我們一般將CPU代碼中互相不依賴的的for循環(huán)適當(dāng)替換成CUDA代碼。

這份代碼打印了8個(gè)數(shù)字，核函數(shù)有一個(gè)參數(shù)N，N = 8，假如我們只想打印5個(gè)數(shù)字呢？當(dāng)前的執(zhí)行配置共2 * 4 = 8個(gè)線程，線程數(shù)8與要執(zhí)行的次數(shù)5不匹配，不過(guò)我們已經(jīng)在代碼里寫好了if (idx < N)的判斷語(yǔ)句，判斷會(huì)幫我們過(guò)濾不需要的計(jì)算。我們只需要把N = 5傳遞給gpu_print函數(shù)中就好，CUDA仍然會(huì)啟動(dòng)8個(gè)thread，但是大于等于N的thread不進(jìn)行計(jì)算。
注意，當(dāng)線程數(shù)與計(jì)算次數(shù)不一致時(shí)，一定要使用這樣的判斷語(yǔ)句，以保證某個(gè)線程的計(jì)算不會(huì)影響其他線程的數(shù)據(jù)。

Block大小設(shè)置

不同的執(zhí)行配置會(huì)影響GPU程序的速度，一般需要多次調(diào)試才能找到較好的執(zhí)行配置，在實(shí)際編程中，執(zhí)行配置[gridDim, blockDim]應(yīng)參考下面的方法：

• block運(yùn)行在SM上，不同硬件架構(gòu)（Turing、Volta、Pascal…）的CUDA核心數(shù)不同，一般需要根據(jù)當(dāng)前硬件來(lái)設(shè)置block的大小blockDim（執(zhí)行配置中第二個(gè)參數(shù)）。一個(gè)block中的thread數(shù)最好是32、128、256的倍數(shù)。注意，限于當(dāng)前硬件的設(shè)計(jì)，block大小不能超過(guò)1024。
• grid的大小gridDim（執(zhí)行配置中第一個(gè)參數(shù)），即一個(gè)grid中block的個(gè)數(shù)可以由總次數(shù)N除以blockDim，并向上取整。

例如，我們想并行啟動(dòng)1000個(gè)thread，可以將blockDim設(shè)置為128，1000 ÷ 128 = 7.8，向上取整為8。使用時(shí)，執(zhí)行配置可以寫成gpuWork8, 128，CUDA共啟動(dòng)8 * 128 = 1024個(gè)thread，實(shí)際計(jì)算時(shí)只使用前1000個(gè)thread，多余的24個(gè)thread不進(jìn)行計(jì)算。

注意，這幾個(gè)變量比較容易混淆，再次明確一下：blockDim是block中thread的個(gè)數(shù)，一個(gè)block中的threadIdx最大不超過(guò)blockDim；gridDim是grid中block的個(gè)數(shù)，一個(gè)grid中的blockIdx最大不超過(guò)gridDim。

以上討論中，block和grid大小均是一維，實(shí)際編程使用的執(zhí)行配置常常更復(fù)雜，block和grid的大小可以設(shè)置為二維甚至三維，如下圖所示。這部分內(nèi)容將在下篇文章中討論。

Thread Block Grid

內(nèi)存分配

前文提到，GPU計(jì)算時(shí)直接從顯存中讀取數(shù)據(jù)，因此每當(dāng)計(jì)算時(shí)要將數(shù)據(jù)從主存拷貝到顯存上，用CUDA的術(shù)語(yǔ)來(lái)說(shuō)就是要把數(shù)據(jù)從主機(jī)端拷貝到設(shè)備端。CUDA強(qiáng)大之處在于它能自動(dòng)將數(shù)據(jù)從主機(jī)和設(shè)備間相互拷貝，不需要程序員在代碼中寫明。這種方法對(duì)編程者來(lái)說(shuō)非常方便，不必對(duì)原有的CPU代碼做大量改動(dòng)。

我們以一個(gè)向量加法為例，編寫一個(gè)向量加法的核函數(shù)如下：

@cuda.jit
def gpu_add(a, b, result, n):
    # a, b為輸入向量，result為輸出向量
    # 所有向量都是n維
    # 得到當(dāng)前thread的索引
    idx = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x
    if idx < n:
        result[idx] = a[idx] + b[idx]

初始化兩個(gè)2千萬(wàn)維的向量，作為參數(shù)傳遞給核函數(shù)：

n = 20000000
x = np.arange(n).astype(np.int32)
y = 2 * x
gpu_result = np.zeros(n)

# CUDA執(zhí)行配置
threads_per_block = 1024
blocks_per_grid = math.ceil(n / threads_per_block)

gpu_add[blocks_per_grid, threads_per_block](x, y, gpu_result, n)

把上述代碼整合起來(lái)，與CPU代碼做對(duì)比，并驗(yàn)證結(jié)果正確性：

from numba import cuda
import numpy as np
import math
from time import time

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

def main():
    n = 20000000
    x = np.arange(n).astype(np.int32)
    y = 2 * x

    gpu_result = np.zeros(n)
    cpu_result = np.zeros(n)

    threads_per_block = 1024
    blocks_per_grid = math.ceil(n / threads_per_block)
    start = time()
    gpu_add[blocks_per_grid, threads_per_block](x, y, gpu_result, n)
    cuda.synchronize()
    print("gpu vector add time " + str(time() - start))
    start = time()
    cpu_result = np.add(x, y)
    print("cpu vector add time " + str(time() - start))

    if (np.array_equal(cpu_result, gpu_result)):
        print("result correct")

if __name__ == "__main__":
    main()

運(yùn)行結(jié)果，GPU代碼竟然比CPU代碼慢10+倍！

gpu vector add time 13.589356184005737
cpu vector add time 1.2823548316955566
result correct

說(shuō)好的GPU比CPU快幾十倍上百倍呢？這里GPU比CPU慢很多原因主要在于：

1. 向量加法的這個(gè)計(jì)算比較簡(jiǎn)單，CPU的numpy已經(jīng)優(yōu)化到了極致，無(wú)法突出GPU的優(yōu)勢(shì)，我們要解決實(shí)際問(wèn)題往往比這個(gè)復(fù)雜得多，當(dāng)解決復(fù)雜問(wèn)題時(shí)，優(yōu)化后的GPU代碼將遠(yuǎn)快于CPU代碼。
2. 這份代碼使用CUDA默認(rèn)的統(tǒng)一內(nèi)存管理機(jī)制，沒(méi)有對(duì)數(shù)據(jù)的拷貝做優(yōu)化。CUDA的統(tǒng)一內(nèi)存系統(tǒng)是當(dāng)GPU運(yùn)行到某塊數(shù)據(jù)發(fā)現(xiàn)不在設(shè)備端時(shí)，再去主機(jī)端中將數(shù)據(jù)拷貝過(guò)來(lái)，當(dāng)執(zhí)行完核函數(shù)后，又將所有的內(nèi)存拷貝回主存。在上面的代碼中，輸入的兩個(gè)向量是只讀的，沒(méi)必要再拷貝回主存。
3. 這份代碼沒(méi)有做流水線優(yōu)化。CUDA并非同時(shí)計(jì)算2千萬(wàn)個(gè)數(shù)據(jù)，一般分批流水線工作：一邊對(duì)2000萬(wàn)中的某批數(shù)據(jù)進(jìn)行計(jì)算，一邊將下一批數(shù)據(jù)從主存拷貝過(guò)來(lái)。計(jì)算占用的是CUDA核心，數(shù)據(jù)拷貝占用的是總線，所需資源不同，互相不存在競(jìng)爭(zhēng)關(guān)系。這種機(jī)制被稱為流水線。這部分內(nèi)容將在下篇文章中討論。

原因2中本該程序員動(dòng)腦思考的問(wèn)題交給了CUDA解決，增加了時(shí)間開銷，所以CUDA非常方便的統(tǒng)一內(nèi)存模型缺點(diǎn)是計(jì)算速度慢。針對(duì)原因2，我們可以繼續(xù)優(yōu)化這個(gè)程序，告知GPU哪些數(shù)據(jù)需要拷貝到設(shè)備，哪些需要拷貝回主機(jī)。

from numba import cuda
import numpy as np
import math
from time import time

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

def main():
    n = 20000000
    x = np.arange(n).astype(np.int32)
    y = 2 * x

    # 拷貝數(shù)據(jù)到設(shè)備端
    x_device = cuda.to_device(x)
    y_device = cuda.to_device(y)
    # 在顯卡設(shè)備上初始化一塊用于存放GPU計(jì)算結(jié)果的空間
    gpu_result = cuda.device_array(n)
    cpu_result = np.empty(n)

    threads_per_block = 1024
    blocks_per_grid = math.ceil(n / threads_per_block)
    start = time()
    gpu_add[blocks_per_grid, threads_per_block](x_device, y_device, gpu_result, n)
    cuda.synchronize()
    print("gpu vector add time " + str(time() - start))
    start = time()
    cpu_result = np.add(x, y)
    print("cpu vector add time " + str(time() - start))

    if (np.array_equal(cpu_result, gpu_result.copy_to_host())):
        print("result correct!")

if __name__ == "__main__":
    main()

這段代碼的運(yùn)行結(jié)果為：

gpu vector add time 0.19940638542175293
cpu vector add time 1.132070541381836
result correct!

至此，可以看到GPU速度終于比CPU快了很多。

Numba對(duì)Numpy的比較友好，編程中一定要使用Numpy的數(shù)據(jù)類型。用到的比較多的內(nèi)存分配函數(shù)有：

• cuda.device_array()： 在設(shè)備上分配一個(gè)空向量，類似于numpy.empty()
• cuda.to_device()：將主機(jī)的數(shù)據(jù)拷貝到設(shè)備

ary = np.arange(10)
device_ary = cuda.to_device(ary)

• cuda.copy_to_host()：將設(shè)備的數(shù)據(jù)拷貝回主機(jī)

host_ary = device_ary.copy_to_host()

總結(jié)

Python Numba庫(kù)可以調(diào)用CUDA進(jìn)行GPU編程，CPU端被稱為主機(jī)，GPU端被稱為設(shè)備，運(yùn)行在GPU上的函數(shù)被稱為核函數(shù)，調(diào)用核函數(shù)時(shí)需要有執(zhí)行配置，以告知CUDA以多大的并行粒度來(lái)計(jì)算。使用GPU編程時(shí)要合理地將數(shù)據(jù)在主機(jī)和設(shè)備間互相拷貝。

GPU程序執(zhí)行流程

CUDA編程的基本流程為：

1. 初始化，并將必要的數(shù)據(jù)拷貝到GPU設(shè)備的顯存上。
2. 使用某個(gè)執(zhí)行配置，以一定的并行粒度調(diào)用CUDA核函數(shù)。
3. CPU和GPU異步計(jì)算。
4. 將GPU計(jì)算結(jié)果拷貝回主機(jī)。

關(guān)于Python學(xué)習(xí)資料：

在學(xué)習(xí)python中有任何困難不懂的可以微信掃描下方CSDN官方認(rèn)證二維碼加入python交流學(xué)習(xí)
多多交流問(wèn)題，互幫互助，這里有不錯(cuò)的學(xué)習(xí)教程和開發(fā)工具。

??[[CSDN大禮包：《python安裝包&全套學(xué)習(xí)資料》免費(fèi)分享]]（安全鏈接，放心點(diǎn)擊）

一、Python所有方向的學(xué)習(xí)路線

Python所有方向的技術(shù)點(diǎn)做的整理，形成各個(gè)領(lǐng)域的知識(shí)點(diǎn)匯總，它的用處就在于，你可以按照上面的知識(shí)點(diǎn)去找對(duì)應(yīng)的學(xué)習(xí)資源，保證自己學(xué)得較為全面。

二、Python必備開發(fā)工具

三、Python視頻合集

觀看零基礎(chǔ)學(xué)習(xí)視頻，看視頻學(xué)習(xí)是最快捷也是最有效果的方式，跟著視頻中老師的思路，從基礎(chǔ)到深入，還是很容易入門的。

四、實(shí)戰(zhàn)案例

光學(xué)理論是沒(méi)用的，要學(xué)會(huì)跟著一起敲，要?jiǎng)邮謱?shí)操，才能將自己的所學(xué)運(yùn)用到實(shí)際當(dāng)中去，這時(shí)候可以搞點(diǎn)實(shí)戰(zhàn)案例來(lái)學(xué)習(xí)。

五、Python練習(xí)題

檢查學(xué)習(xí)結(jié)果。

六、面試資料

我們學(xué)習(xí)Python必然是為了找到高薪的工作，下面這些面試題是來(lái)自阿里、騰訊、字節(jié)等一線互聯(lián)網(wǎng)大廠最新的面試資料，并且有阿里大佬給出了權(quán)威的解答，刷完這一套面試資料相信大家都能找到滿意的工作。


最后，千萬(wàn)別辜負(fù)自己當(dāng)時(shí)開始的一腔熱血，一起變強(qiáng)大變優(yōu)秀。文章來(lái)源地址http://www.zghlxwxcb.cn/news/detail-434766.html

到了這里，關(guān)于GPU加速02:超詳細(xì)Python Cuda零基礎(chǔ)入門教程，沒(méi)有顯卡也能學(xué)！的文章就介紹完了。如果您還想了解更多內(nèi)容，請(qǐng)?jiān)谟疑辖撬阉鱐OY模板網(wǎng)以前的文章或繼續(xù)瀏覽下面的相關(guān)文章，希望大家以后多多支持TOY模板網(wǎng)！