目錄¹

介紹篇
執行緒篇
程序篇
非同步篇
GPU篇
分散式篇

準備

1. 需要有支援CUDA的Nvidia顯示卡
   linux檢視顯示卡資訊：lspci | grep -i vga
   使用nvidia顯示卡可以這樣檢視：lspci | grep -i nvidia
   上一個命令可以得到類似"03.00.0"的顯示卡代號，檢視詳細資訊：lspic -v -s 03.00.0
   檢視顯示卡使用情況（nvidia專用）：nvidia-smi
   持續週期性輸出使用情況（1秒1次）：watch -n 1 nvidia-smi
2. 需要安裝pycuda（linux安裝：apt install python3-pycuda）

基本使用

1.底層操作

0. 準備工作: sudo apt install python3-pycuda
1. 在cpu使用的記憶體中建立矩陣
2. 將矩陣從cpu使用的記憶體移動到gpu的視訊記憶體中
3. 編輯c語言程式碼，讓gpu計算
4. 將矩陣從gpu視訊記憶體移動到cpu使用的記憶體中

核心與執行緒層級：
CUDA程式的一個最重要元素就是核心（kernel），它代表可以並行執行的程式碼
每個核心的執行均有叫做執行緒（thread）的計算單元完成，與cpu的執行緒不同，gpu執行緒更加輕量，上下文切換不會影響效能
為了確定執行一個核心所需的執行緒數機器邏輯組織形式，CUDA定義了一個二層結構。在最高一層，定義了所謂的區塊網格（grid of blocks），這個網格代表了執行緒區塊所在的二維結構，而這些執行緒區塊則是三維的（簡單來說，一個cuda結構包含多個blocks，每個blocks包含多個thread）（在下面還會對每個執行緒區塊細分操作）

一個執行緒區塊會被指派給一個流式多處理器(SM)，然後這些執行緒被進一步劃分為被稱為warp的執行緒組，其大小有GPU的架構決定
為了充分發揮SM本身的併發性，同一組內的執行緒必須執行相同的指令，否則會出現執行緒分歧（divergence of thread）

示例：
用gpu將矩陣每個元素x2

import pycuda.driver as cuda
import pycuda.autoinit  # init GPU
from pycuda.compiler import SourceModule

import numpy as np

# 1.cpu create matrix
 

a = np.random.randn(5, 5)	# matrix:m*n
a = a.astype(np.float32)	# nvidia only support float calculate

# 2.move to gpu from cpu
a_gpu = cuda.mem_alloc(a.nbytes)	# alloc memory of gpu, this is 1 dim
cuda.memcpy_htod(a_gpu, a)		# copy cpu memory to gpu memory

# 3.gpu calculate
# create module of gpu calculate by c
mod = SourceModule('''
	__global__ void doubleMatrix(float *a)
	{
		int idx = threadIdx.x + threadIdx.y * 5;	// (x,y,z), gpu -> sm -> warp
		a[idx] *= 2;
	}
''')

func = mod.get_function('doubleMatrix')	# get function from module
func(a_gpu, block = (5, 5, 1))		# set var and thread number from (x,y,z) orient to function

# 4.move to cpu from gpu
a_doubled = np.empty_like(a)		# create memory of cpu
cuda.memcpy_dtoh(a_doubled, a_gpu)	# copy gpu memory to cpu memory

print('original matrix')
print(a)
print('double matrix')
print(a_doubled)

注1：import pycuda.autoinit 語句自動根據GPU可用性和數量選擇要使用的GPU，這將建立一個接下來的程式碼執行中所需的GPU上下文（只需匯入即可完成）

注: astype(numpy.float32)：將矩陣中的項轉換為單精度模式，因為許多Nvidia顯示卡只支援單精度

注3：在呼叫gpu的c函式時，通過block引數設定，分配執行緒的方式，（5, 5, 1）是對應這gpu的(x, y, x)的分配
在c函式中，threadIdx是一個結構體，它有三個欄位x、y、z，每個執行緒中的這個變數都不同（結合線程層級理解），故用此索引陣列，由於動態分配的gpu記憶體是一維陣列，所以需要在c函式內，使用threadIdx.y乘以矩陣每行的元素個數轉換
更多詳情：自行搜尋cuda的執行緒區塊劃分

注4：gpu執行的c函式中 __global__ 關鍵字表示該函式是一個核心函式，必須從主機上呼叫才能在gpu裝置上生成執行緒層級
涉及到pycuda記憶體，為了最大限度地利用可用資源，在支援CUDA的GPU顯示卡中，有4類記憶體：
暫存器(registers)：每個執行緒將被分配一個暫存器，每個執行緒只能訪問自身的暫存器，即使同屬於一個執行緒區塊
共享儲存器(shared memory)：在共享儲存器中，每個執行緒區塊都有一個其內部執行緒共享的記憶體，這部分記憶體速度極快
常數儲存器(constant memory)：一個網格中的所有執行緒一直都可以訪問這部分記憶體，但只能在讀取時訪問。常數儲存器中的資料在應用持續期間一直存在
全域性儲存器(global memory)，：所有網格中的執行緒（也包含所有核心）都可訪問全域性儲存器
更多詳情：自行搜尋PyCUDA記憶體模型

2.python封裝控制

用gpuarray呼叫核心，它可以直接將資料儲存在計算裝置(gpu)中，並在該裝置中進行計算

示例：
用gpu將矩陣每個元素x2

import pycuda.autoinit
import pycuda.gpuarray as gpuarray

import numpy as np

x = np.random.randn(5, 5)
x_gpu = gpuarray.to_gpu(x)

x_doubled = (2 * x_gpu).get()

print('x:')
print(x)
print('x_doubled:')
print(x_doubled)

其它

還有其它的庫可以方便對cuda程式設計，如NumbaPro是一個Python編譯器，提供了基於CUDA的API程式設計介面，可以編寫CUDA持續，它專門設計用來執行與陣列相關的計算任務，和廣泛使用的numpy庫類似
NumbaPro: 對GPU程式設計的庫，提供許多的數值計算庫，GPU加速庫