待學習：
- Synchronized
- Overlaped
英偉達CUDA介紹
CUDA6中的Unified memory
CUDA簡介
CPU & GPU
GPU包括更多的運算核心，其特別適合數據並行的計算密集型任務，如大型矩陣運算，而CPU的運算核心較少，但是其可以實現復雜的邏輯運算，因此其適合控制密集型任務。另外，CPU上的線程是重量級的，上下文切換開銷大，但是GPU由於存在很多核心，其線程是輕量級的。因此，基於CPU+GPU的異構計算平台可以優勢互補，CPU負責處理邏輯復雜的串行程序，而GPU重點處理數據密集型的並行計算程序，從而發揮最大功效。
CUDA是NVIDIA公司所開發的GPU編程模型，它提供了GPU編程的簡易接口，基於CUDA編程可以構建基於GPU計算的應用程序
CUDA編程
CUDA編程模型是一個異構模型，需要CPU和GPU協同工作。在CUDA中，host和device是兩個重要的概念，我們用host指代CPU及其內存，而用device指代GPU及其內存。CUDA程序中既包含host程序，又包含device程序，它們分別在CPU和GPU上運行。同時，host與device之間可以進行通信，這樣它們之間可以進行數據拷貝。典型的CUDA程序的執行流程如下：
- 分配host內存，並進行數據初始化；
- 分配device內存，並從host將數據拷貝到device上；
- 調用CUDA的核函數在device上完成指定的運算；
- 將device上的運算結果拷貝到host上；
- 釋放device和host上分配的內存。
CUDA邏輯層
kernel是CUDA中一個重要的概念，kernel是在device上線程中並行執行的函數，核函數用__global__符號聲明，在調用時需要用<<<grid, block>>>來指定kernel要執行的線程數量，在CUDA中，每一個線程都要執行核函數，並且每個線程會分配一個唯一的線程號thread ID。
kernel在device上執行時實際上是啟動很多線程，一個kernel所啟動的所有線程稱為一個網格（grid），同一個網格上的線程共享相同的全局內存空間，grid是線程結構的第一層次，而網格又可以分為很多線程塊（block），一個線程塊里面包含很多線程，這是第二個層次。
CUDA物理層
一個kernel實際上會啟動很多線程，這些線程是邏輯上並行的，但是在物理層卻並不一定。這其實和CPU的多線程有類似之處，多線程如果沒有多核支持，在物理層也是無法實現並行的。但是好在GPU存在很多CUDA核心，充分利用CUDA核心可以充分發揮GPU的並行計算能力。
GPU硬件的一個核心組件是SM，前面已經說過，SM是英文名是 Streaming Multiprocessor，翻譯過來就是流式多處理器。SM的核心組件包括CUDA核心，共享內存，寄存器等，SM可以並發地執行數百個線程，並發能力就取決於SM所擁有的資源數。當一個kernel被執行時，它的gird中的線程塊被分配到SM上，一個線程塊只能在一個SM上被調度。SM一般可以調度多個線程塊，這要看SM本身的能力。那么有可能一個kernel的各個線程塊被分配多個SM，所以grid只是邏輯層，而SM才是執行的物理層。
SM采用的是SIMT (Single-Instruction, Multiple-Thread，單指令多線程)架構，基本的執行單元是線程束（wraps)，線程束包含32個線程，這些線程同時執行相同的指令，但是每個線程都包含自己的指令地址計數器和寄存器狀態，也有自己獨立的執行路徑。所以盡管線程束中的線程同時從同一程序地址執行，但是可能具有不同的行為，比如遇到了分支結構，一些線程可能進入這個分支，但是另外一些有可能不執行，它們只能死等，因為GPU規定線程束中所有線程在同一周期執行相同的指令，線程束分化會導致性能下降。
當線程塊被划分到某個SM上時，它將進一步划分為多個線程束，因為這才是SM的基本執行單元，但是一個SM同時並發的線程束數是有限的。這是因為資源限制，SM要為每個線程塊分配共享內存，而也要為每個線程束中的線程分配獨立的寄存器。所以SM的配置會影響其所支持的線程塊和線程束並發數量。總之，就是網格和線程塊只是邏輯划分，一個kernel的所有線程其實在物理層是不一定同時並發的。所以kernel的grid和block的配置不同，性能會出現差異，這點是要特別注意的。
還有，由於SM的基本執行單元是包含32個線程的線程束，所以block大小一般要設置為32的倍數。
CPU和GPU間數據傳輸
Pinned Host Memory
參考
主機(CPU)數據分配默認是分頁的。GPU不能直接訪問分頁主機內存中的數據。
CUDA驅動先分配(allocate)一個臨時的page-locked,or pinned host array，再將主機的數據拷貝到pinned array，再轉換到device memory。
我們可以直接將主機數據分配到pinned memory中，減少一次傳輸。在C++中使用cudaMallocHost()函數。
Unified Memory in CUDA 6
參考自英偉達官方文檔
CPU和GPU中的數據必須分開放置在各自的存儲中，通過程序進行復制傳遞。
Unified Memory創建了一個pool of managed memory來共享二者之間的數據。只通過同一個指針，GPU和CPU就都可以訪問這塊managed memory。
實際上，系統會自動在host和device之間遷移管理managed memory中的數據，使得數據看起來像是在一塊可共享的內存中。所以我們只需要進行一次內存分配，就可以得到一個能從host和device訪問的指向數據的指針。
英偉達CUDA介紹
CUDA簡介
CPU & GPU
GPU包括更多的運算核心，其特別適合數據並行的計算密集型任務，如大型矩陣運算，而CPU的運算核心較少，但是其可以實現復雜的邏輯運算，因此其適合控制密集型任務。另外，CPU上的線程是重量級的，上下文切換開銷大，但是GPU由於存在很多核心，其線程是輕量級的。因此，基於CPU+GPU的異構計算平台可以優勢互補，CPU負責處理邏輯復雜的串行程序，而GPU重點處理數據密集型的並行計算程序，從而發揮最大功效。
CUDA是NVIDIA公司所開發的GPU編程模型，它提供了GPU編程的簡易接口，基於CUDA編程可以構建基於GPU計算的應用程序
CUDA編程
CUDA編程模型是一個異構模型，需要CPU和GPU協同工作。在CUDA中，host和device是兩個重要的概念，我們用host指代CPU及其內存，而用device指代GPU及其內存。CUDA程序中既包含host程序，又包含device程序，它們分別在CPU和GPU上運行。同時，host與device之間可以進行通信，這樣它們之間可以進行數據拷貝。典型的CUDA程序的執行流程如下：
- 分配host內存，並進行數據初始化；
- 分配device內存，並從host將數據拷貝到device上；
- 調用CUDA的核函數在device上完成指定的運算；
- 將device上的運算結果拷貝到host上；
- 釋放device和host上分配的內存。
CUDA邏輯層
kernel是CUDA中一個重要的概念，kernel是在device上線程中並行執行的函數，核函數用__global__符號聲明，在調用時需要用<<<grid, block>>>來指定kernel要執行的線程數量，在CUDA中，每一個線程都要執行核函數，並且每個線程會分配一個唯一的線程號thread ID。
kernel在device上執行時實際上是啟動很多線程，一個kernel所啟動的所有線程稱為一個網格（grid），同一個網格上的線程共享相同的全局內存空間，grid是線程結構的第一層次，而網格又可以分為很多線程塊（block），一個線程塊里面包含很多線程，這是第二個層次。
CUDA物理層
一個kernel實際上會啟動很多線程，這些線程是邏輯上並行的，但是在物理層卻並不一定。這其實和CPU的多線程有類似之處，多線程如果沒有多核支持，在物理層也是無法實現並行的。但是好在GPU存在很多CUDA核心，充分利用CUDA核心可以充分發揮GPU的並行計算能力。
GPU硬件的一個核心組件是SM，前面已經說過，SM是英文名是 Streaming Multiprocessor，翻譯過來就是流式多處理器。SM的核心組件包括CUDA核心，共享內存，寄存器等，SM可以並發地執行數百個線程，並發能力就取決於SM所擁有的資源數。當一個kernel被執行時，它的gird中的線程塊被分配到SM上，一個線程塊只能在一個SM上被調度。SM一般可以調度多個線程塊，這要看SM本身的能力。那么有可能一個kernel的各個線程塊被分配多個SM，所以grid只是邏輯層，而SM才是執行的物理層。
SM采用的是SIMT (Single-Instruction, Multiple-Thread，單指令多線程)架構，基本的執行單元是線程束（wraps)，線程束包含32個線程，這些線程同時執行相同的指令，但是每個線程都包含自己的指令地址計數器和寄存器狀態，也有自己獨立的執行路徑。所以盡管線程束中的線程同時從同一程序地址執行，但是可能具有不同的行為，比如遇到了分支結構，一些線程可能進入這個分支，但是另外一些有可能不執行，它們只能死等，因為GPU規定線程束中所有線程在同一周期執行相同的指令，線程束分化會導致性能下降。
當線程塊被划分到某個SM上時，它將進一步划分為多個線程束，因為這才是SM的基本執行單元，但是一個SM同時並發的線程束數是有限的。這是因為資源限制，SM要為每個線程塊分配共享內存，而也要為每個線程束中的線程分配獨立的寄存器。所以SM的配置會影響其所支持的線程塊和線程束並發數量。總之，就是網格和線程塊只是邏輯划分，一個kernel的所有線程其實在物理層是不一定同時並發的。所以kernel的grid和block的配置不同，性能會出現差異，這點是要特別注意的。
還有，由於SM的基本執行單元是包含32個線程的線程束，所以block大小一般要設置為32的倍數。
CPU和GPU間數據傳輸
Pinned Host Memory
參考
主機(CPU)數據分配默認是分頁的。GPU不能直接訪問分頁主機內存中的數據。
CUDA驅動先分配(allocate)一個臨時的page-locked,or pinned host array，再將主機的數據拷貝到pinned array，再轉換到device memory。
我們可以直接將主機數據分配到pinned memory中，減少一次傳輸。在C++中使用cudaMallocHost()函數。
Unified Memory in CUDA 6
參考自英偉達官方文檔
CPU和GPU中的數據必須分開放置在各自的存儲中，通過程序進行復制傳遞。
Unified Memory創建了一個pool of managed memory來共享二者之間的數據。只通過同一個指針，GPU和CPU就都可以訪問這塊managed memory。
實際上，系統會自動在host和device之間遷移管理managed memory中的數據，使得數據看起來像是在一塊可共享的內存中。所以我們只需要進行一次內存分配，就可以得到一個能從host和device訪問的指向數據的指針。
庫函數
cudaDeviceSynchronize()：用在host code中，讓CPU等待直到所有GPU操作結束。cudaMemcpy()：只能用於默認的stream0中，自帶syn功能，會阻塞調用它的CPU線程直到拷貝完成。cudaMemcpyAsync()：可以接受一個stream參數，當作用在==非默認stream0且操作內存為pinned memory時，可以在拷貝完成前將控制權交還給調用線程。
Stream
只有不同stream中的CUDA操作才能並發交錯執行。
當沒有設置stream時使用默認的stream，此時host和device完全同步，就相當於在每個cuda操作后面加了一個cudaDeviceSynchronize()
參考資料：英偉達官方PPT，Stream and Concurrency
A sequence of operations that execute in issue-order on GPU
CUDA operations in different streams may run concurrently
CUDA operations from different streams may be interleaved
並發的條件：
- CUDA操作必須在不同的，非零的stream上
- cudaMemcpyAsync必須作用在
pinned memory上(用cudaMallocHost()或者cudaHostAlloc()分配的內存) - 有足夠的資源：
- 不同方向的cudaMemcpyAsyncs
- 設備資源(SMEM, registers, blocks等)
當沒有設置stream時使用默認的stream，此時host和device完全同步，就相當於在每個cuda操作后面加了一個cudaDeviceSynchronize()（該函數作用？？？）
