您现在的位置：程式師世界 >> 編程語言 > >> 更多編程語言 >> Python

Python CUDA 編程 - 3 - GPU編程介紹

編輯：Python

本文介紹GPU編程的一些重要概念。

GPU編程

GPU編程與CPU編程的思考角度不盡相同，舉皮皮魯老師的一個例子：

以加法計算為例，CPU就像大學數學教授，GPU就像幾千個小學生，現在需要不借助外界，只通過紙筆，對2000個數字進行加法計算，得到1000個加法結果，在這個過程中，大學教授要協調指揮小學生完成任務。

在計算過程中，每個小學生需要按照大學教授的提出的規范，基於一個加法函數，完成計算。每個小學生就像GPU的一個計算核心，加法函數就是核函數，一個小學生完成本次計算就像一次線程計算。在整個計算過程中，只能通過紙筆交流，無論是計算任務本身，還是計算的中間結果都需要落地到紙上進行計算，作為記錄的紙就像是計算機中的存儲，

假設我們有2000個數字需要加到一起，得到1000個加法結果。如果只依賴20個大學教授，那每個教授需要執行50次計算，耗時很長。如果大學教授可以借助1000個小學生同時計算，那麼大學教授可以這樣設計GPU並行算法：

設計一個加法函數，加法函數可以將兩個數字加在一起。
每個小學生分配2個數字，使用加法函數，對這2個數字執行計算。
大學教授給1000個小學生分配數字，並告知他們使用怎樣的加法函數進行計算。
1000個小學生同時計算，得到1000個結果數字，寫到紙上，返回給大學教授。

實際上，CUDA並行算法和上面的流程基本相似，就是設計核函數，在存儲上合理分配數據，告知GPU以一定的並行度執行配置來並行計算。核函數的設計與所要解決的問題本身高度相關。

主機與設備

在CUDA中，CPU和主存被稱為主機（Host），GPU和顯存（顯卡內存）被稱為設備（Device），CPU無法直接讀取顯存數據，GPU無法直接讀取主存數據，主機與設備必須通過總線（Bus）相互通信。

GPU程序與CPU程序的區別

CPU程序

初始化。
CPU計算。
得到計算結果。

GPU程序

初始化，並將必要的數據拷貝到GPU設備的顯存上。
CPU調用GPU函數，啟動GPU多個核心同時進行計算。
CPU與GPU異步計算。
將GPU計算結果拷貝回主機端，得到計算結果

Thread層次結構

在進行GPU並行編程時，需要定義執行配置來告知以怎樣的方式去並行執行核函數。CUDA將核函數所定義的運算稱為線程（Thread），多個線程組成一個塊（Block），多個塊組成網格（Grid）。這樣一個Grid可以定義成千上萬個線程，也就解決了並行執行上萬次操作的問題。

實際上，線程（Thread）是一個編程上的軟件概念。從硬件來看，Thread運行在一個CUDA核心上，多個Thread組成的Block運行在Streaming Multiprocessor（SM），多個Block組成的Grid運行在一個GPU顯卡上。

)

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

某個Thread在整個Grid中的位置編號為：threadIdx.x + blockIdx.x * blockDim.x。

Block大小設置

不同的執行配置會影響GPU程序的速度，一般需要多次調試才能找到較好的執行配置，在實際編程中，執行配置[gridDim, blockDim]應參考下面的方法：

Block運行在SM上，不同硬件架構（Turing、Volta、Pascal…）的CUDA核心數不同，一般需要根據當前硬件來設置Block的大小blockDim（執行配置中第二個參數）。一個Block中的Thread數最好是32、128、256的倍數。注意，限於當前硬件的設計，Block大小不能超過1024。
Grid的大小gridDim（執行配置中第一個參數），即一個Grid中Block的個數可以由總次數N除以blockDim，並向上取整。

例如，我們想並行啟動1000個Thread，可以將blockDim設置為128，1000 ÷ 128 = 7.8，向上取整為8。使用時，執行配置可以寫成gpuWork[8, 128]()，CUDA共啟動8 * 128 = 1024個Thread，實際計算時只使用前1000個Thread，多余的24個Thread不進行計算。

這幾個變量比較容易混淆，再次明確一下：blockDim是Block中Thread的個數，一個Block中的threadIdx最大不超過blockDim；gridDim是Grid中Block的個數，一個Grid中的blockIdx最大不超過gridDim。

以上討論中，Block和Grid大小均是一維，實際編程使用的執行配置常常更復雜，Block和Grid的大小可以設置為二維甚至三維：

一個二維的執行配置如上圖所示，其中，每個Block有(3 * 4)個Thread，每個Grid有(2 * 3)個Block。二維塊大小為 (Dx, Dy)，某個線程號 (x, y) 的公式為 (x + y Dx)；三維塊大小為 (Dx, Dy, Dz)，某個線程號*(x, y, z)* 的公式為 (x + y Dx + z Dx Dy)。各個內置變量中.x.y和.z為不同維度下的值。

例如，一個二維配置，某個線程在矩陣中的位置可以表示為：

col = cuda.threadIdx.y + cuda.blockDim.y * cuda.blockIdx.y
row = cuda.threadIdx.x + cuda.blockDim.x * cuda.blockIdx.x

如何將二維Block映射到自己的數據上並沒有固定的映射方法，一般情況將.x映射為矩陣的行，將.y映射為矩陣的列。Numba提供了一個更簡單的方法幫我們計算線程的編號：

row, col = cuda.grid(2)

其中，參數2表示這是一個2維的執行配置。1維或3維的時候，可以將參數改為1或3。

對應的執行配置也要改為二維：

threads_per_block = (16, 16)
blocks_per_grid = (32, 32)
gpu_kernel[blocks_per_grid, threads_per_block]

(16, 16)的二維Block是一個常用的配置，共256個線程。之前也曾提到過，每個Block的Thread個數最好是128、256或512，這與GPU的硬件架構高度相關。

內存分配

前文提到，GPU計算時直接從顯存中讀取數據，因此每當計算時要將數據從主存拷貝到顯存上，用CUDA的術語來說就是要把數據從主機端拷貝到設備端。用小學生計算的例子來解釋，大學教授需要將計算任務寫在紙上，分發給各組小學生。CUDA強大之處在於它能自動將數據從主機和設備間相互拷貝，不需要程序員在代碼中寫明。這種方法對編程者來說非常方便，不必對原有的CPU代碼做大量改動。