? Numba是一個針對Python的開源JIT編譯器，由Anaconda公司主導開發(fā)，可以對Python原生代碼進行CPU和GPU加速。Numba對NumPy數(shù)組和函數(shù)非常友好。

? ? ? ? 使用Numba非常方便，只需要在Python原生函數(shù)上增加一個裝飾器（Decorator）。Numba會將這些函數(shù)使用即時編譯JIT方式編譯成機器碼，這些代碼將以近乎機器碼的速度運行。

2.GPU加速：?

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

使用from numba import cuda引入cuda庫

在GPU函數(shù)上添加@cuda.jit裝飾符，表示該函數(shù)是一個在GPU設備上運行的函數(shù)，GPU函數(shù)又被稱為核函數(shù)。

主函數(shù)調(diào)用GPU核函數(shù)時，需要添加如[1, 2]這樣的執(zhí)行配置，這個配置是在告知GPU以多大的并行粒度同時進行計算。gpu_func[1, 2]()表示開啟一個block塊，每個block塊開啟2個線程并行地執(zhí)行gpu_func函數(shù)，函數(shù)將被并行地執(zhí)行2次。

GPU核函數(shù)的啟動方式是異步的：啟動GPU函數(shù)后，CPU不會等待GPU函數(shù)執(zhí)行完畢才執(zhí)行下一行代碼。必要時，需要調(diào)用cuda.synchronize()，告知CPU等待GPU執(zhí)行完核函數(shù)后，再進行CPU端后續(xù)計算。這個過程被稱為同步。

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

? CUDA提供了一系列內(nèi)置變量，以記錄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。

圖7：GPU的內(nèi)存模型。GPU里面的內(nèi)存分為三種：per thread local memory， per block shared memory，和global memory。在實際編程的時候，需要考慮多用shared memory，少用global memory，因為shared比global的訪問和存取速度快很多。

3 GPU的Hello world: 矩陣相乘

能夠讓人理解GPU編程的 Hello world 程序，便是矩陣相乘。這個紐約大學的教程非常棒，詳細講解了如何編寫GPU程序進行矩陣相乘。我當時學習Numba和CUDA，這個教程發(fā)揮了很大的作用。

3.1介紹幾個名詞

首先，要弄清楚下面6個名詞的概念。編寫GPU程序，其實是一種CPU和GPU之間的“互動”，所謂的異構開發(fā)。

Host。代表CPU。

Device。代表GPU。

Host memory。RAM內(nèi)存。

Device memory。GPU上的存儲。

Kernal function。GPU函數(shù)，執(zhí)行在device上面，調(diào)用者是host。

Device function。GPU函數(shù)，執(zhí)行在device上面，調(diào)用者是kernal function或者device function

3.2 GPU程序的執(zhí)行流程

圖5可視化了GPU程序的流程：

把數(shù)組的數(shù)據(jù)(numpy.ndarray)從host memory拷貝到device memory上。

配置kernal function的參數(shù)，調(diào)用kernal function。參數(shù)有兩種：一種用中括號[ ]，一種用小括號( )。中括號的參數(shù)是threadperblock和blockpergrid，參考圖6。小括號就是那種普通函數(shù)的參數(shù)。

幾千個線程同時調(diào)用同一個kernal function，在GPU里面進行計算。kernal function的編寫，是一門技術。

把GPU里面的運算結果，從device memory拷貝回host memory。THE END。

3.3 矩陣相乘源代碼

這份Python代碼，有6個函數(shù)，進行 C = A × B \mathbf{C} = \mathbf{A} \times \mathbf{B}C=A×B

cpu_mat_mul(A, B, C), 基于CPU的矩陣相乘，O ( n 3 ) O(n^3)O(n?

3

?)。

cpu_mat_mul_jit(A, B, C)，和cpu_mat_mul相比，唯一的區(qū)別就是加了裝飾器 @jit，這是一種編譯優(yōu)化工具，可以加快幾十倍。

mat_mul_naive_kernal(A, B, C)，矩陣相乘的GPU函數(shù)，使用global memory。

mat_mul_shared_kernal(A, B, C)，矩陣相乘的GPU，使用了shared memory。shared memory比global memory訪問速度快不少。

host_naive(A, B, C)，mat_mul_naive_kernal的“配套設施”，kernal不能獨立存在的，前前后后需要一些代碼輔助。

host_optimized(A, B, C)，mat_mul_shared_kernal的“配套設施”。如果host_naive比cpu_mat_mul快三千倍的話，host_optimized可能快三千五百倍。

Numba also exposes three kinds of GPU memory:

• global device memory

• shared memory

• local memory

Find ways to parallelize sequential code

Minimize data transfers between the host and the device

Adjust kernel launch configuration to maximize device utilization

Ensure global memory accesses are coalesced

Minimize redundant accesses to global memory whenever possible

Avoid different execution paths within the same warp