最近想用cuda來加速三維重建的算法，就先入門了一下cuda。
CUDA C 編程
cuda c時對c/c++進行拓展后形成的變種，兼容c/c++語法，文件類型為'.cu'，編譯器為nvcc。cuda c允許用內核函數來擴展c，調用時由N個不同的線程共執行N次。塊內的線程可以通過共享存儲器共享數據並通過它們的執行力來協調存儲器訪問，aka 通過調用__syncthreads()內部函數來指定內核中的同步點。
相比傳統的cpp，添加了這么幾個方面：
- 函數類型限定符
- 執行配置運算符
- 五個內置變量
- 變量類型限定符
- 其他的還有數學函數，原子函數，紋理讀取，綁定函數等。
函數類型限定符
用來確定時再cpu還是gpu運行，以及這個函數是從cpu還是gpu調用。
- device表示從gpu調用，再gpu運行
- global表示從cpu調用，在gpu執行，也稱kernel函數。
- host表示在cpu上調用，在cpu上執行。
執行配置運算符
用來傳遞核函數的執行參數。
使用__global__聲明說明符定義內核，用<<< ... >>>來為內核指定cuda線程數。每一個線程都有一個唯一的ID，可以通過內置threadIdx來訪問。
<<< ... >>>中可以時int或者dim3類型。
五個內置變量
gridDim;
blockDim;
blockIdx;
threadIdx;
warpSize;
變量類型限定符
__device__; // 表示位於全局內存空間，默認類型
__share__; // 表示位於共享內存空間
__constant__; // 常量內存空間
texture; // 其綁定的變量可以被紋理緩存加速訪問
some unit
- thread：一個cuda的並行程序會被許多個threads來執行。
- block：數個threads會被群組成要給block，同一個block中的threads可以同步，也可以通過shared memory通信。
- grid：多個blocks再構成grid
- warp：GPU執行程序時的調度單位，目前cuda的warp的大小為32，同在一個warp的線程，以不同數據資源執行相同的指令，就是所謂的SIMT。
關於內存：
- 每個線程都有私有本地內存
- 每個線程塊都具有對塊的所有線程可見的共享內存,並且和塊有相同的生存周期
- 所有線程都可以訪問相同的全局內存
所有線程都可以訪問的額外兩個只讀存儲空間：常量內存和紋理內存.
cuda編程還假設主機和設備都在DRAM中保持他們自己的獨立存儲空間，分別成為主機存儲器（host memory)和設備存儲器(device memory)。
統一內存(Unified memory)和托管內存(managed memory)以橋接主機和設備的內存空間。可以從系統的所有CPU和GPU訪問托管內存。
Device Memory
核函數在設備內存之外運行，因此runtime提供分配，釋放和復制內存的功能，以及在主機內存和設備內存之間傳輸數據的功能。
設備存儲器可以分為線性內存(linear memory)或cuda陣列(cuda arrays)。
cuda數組時不透明的內存布局，針對紋理提取進行了優化。
線性內存通常通過cudaMalloc()分配，並使用cudaFree()釋放，主機存儲器和設備存儲器之間的數據傳輸通常使用cudaMemcpy()完成。
還有一些函數：cudaMallocPitch(), cudaMalloc3D(), cudaMemcpy2D(), cudaMemcpy3d().
Shared Memory
共享內存是使用__shared__內存空間說明符分配的。
共享內存一般比全局內存快得多，因此用改利用共享內存替換訪問全局內存的任何機會。
這一塊還不太懂，回頭再摸。