這一部分打算從頭記錄一下CUDA的編程方法和一些物理架構上的特點；從硬件入手，寫一下包括線程束的划分、流水線的調度等等微結構的問題，以及這些物理設備是如何與軟件對應的。下一部分會寫一下cuda中的幾種內存划分，進行數據同步，以及優化cuda運行效率的幾種方法。（傳送門）
1 硬件架構
1.1 Tesla ： G80
不同廠家、版本的GPU內容差別可能會比較大，因此挑出幾款比較經典的GPU，寫一些通用的部分。
GPU最重要的一點是可以並行的實現數據處理。這一點在數據量大、運算復雜度不高的條件下極為適用。可以簡單地把一塊GPU想象成一個超多核的CPU運算部件。這些CPU有自己的寄存器，還有供數據交換用的共享內存、緩存，同時周圍還有取指部件和相應的調度機制，保證指令能夠在之上執行。
G80 在 2006年發布，是Telsa架構下的一塊早期顯卡，擁有現在GPU中一些非常重要的特點，比如SM/SP的划分、多級內存結構、線程束等等，在后文會一一展開。
一塊G80內，設計了多個SM（stream Multiprocessor： 流多處理器）， 每個SM內又包括了多個SP（streaming processor）。而SP正是實現算數功能的核心部件，可以類比CPU之中的ALU單元，只不過其計算能力要差很多。G80中計算單元甚至只有24位，做32位整形數運算的時候需要使用算法模擬。
上圖為一塊G80的架構。L1、L2是兩級緩存，還有一些用於頂點處理等圖形功能的部件；其中拿出這樣的一條來，即為一塊SM：
可以看到，一塊G80上面設計了16塊SM， 其中每個SM上又有8個SP，一共是128個SP。
除了這些計算單元之外，處理器上還有許多其他的部件。每個SM內部的SP之間，可以共享一塊shared memory， 以及一塊指令緩存用於存放指令、一塊常量緩存（c-cache）用來存放常量數據，兩個SFU（特殊運算單元，special function unit）用來做三角函數等較復雜運算，MT issue用來實現多線程下的取指，以及DP（Double Precision Unit）用來做雙精度數。 除去一些運算單元之外，最重要的就是c-cache與shared memory（在G80上為16KB）兩塊數據存儲區。注意這兩個位置的數據只能由SM內部的SP進行訪問，這也是將SP划分在多個SM內的意義所在，后面還會進一步展開。另外注意，如同cpu邊上的cache一樣，cache被故意設計成對程序員透明的，並不能主動控制一個數據到底有沒有存放在緩存里面。(當然，隨着NVCC編譯技術發展，現在比較新的架構上已經允許對L1和L2進行部分的控制，這一點如果之后有時間會展開說一說)
當然，SM之間也有用於數據交換的區域。最主要的是global memory，可以供所有的SM共享，同時有通向CPU上內存的數據總線，which means 通過這里可以從cpu那里獲得數據。容量更大，訪存速度也更慢。
1.2 CPU & GPU
初學CUDA的人可能會對多種不同層次的內存感到困惑。個人理解，多核處理器中一個很重要的問題就是數據的統一。馮諾依曼結構中規定了多級速度不同的存儲器部件，使得訪存可以更快更高效的進行，但多核處理器中還需要維護不同數據間的一致性，不然就會帶來臟讀等問題。這樣一來，把線程划分到不同的層次，有的變量歸為線程私有，有的變量歸為整個block上共享，有的可以全局共享，就可以很方便的提升算法的效率。幾種數據並不互通，當用戶需要的時候可以主動的將一種變量復制到另一個區域，也有利於減少不一致性。
GPU和一個多核cpu很像，主要區別除了存儲部件，還在於指令的調度方式，會在后文線程部分展開。
1.3 other GPU architect
Fermi
fermi 是NVIDIA推出的另一款經典架構，從整體結構上看仍遵從SM--Sp的划分，增加了一些SP數量。注意下圖中的是一塊SM，從前的SM上只有8塊Sp（G80、GT200），而現在增加到了32個；特殊運算單元、緩存的數量也有所增加。
此外還設計了16個LD/ST單元。LD/ST意思就是load & store， 每次可以同時為block內的16個線程（block~線程的關系后面展開）計算目標地址和源地址，加快訪存速度。
SP也有了新的名字--core。其實Core內部已經有不小的變化，每一個Core內部都有一個全流水的整形單元和一個浮點數運算單元，對計算模式下的GPU運行效率有很大改進。在Fermi中，ALU和FPU也支持了32位的運算操作；同時提出了dual scheduler 的設計，兩個線程束可以並行計算。
雙精度數的計算效率提升很大：
Maxwell
maxwell比fermi更晚一些，在2014年發布，其中的一些重要變化如 Core、LD\ST都已經在Fermi中介紹過了。
比較大的變化是Maxwell里面提出了SMM（ Maxwell Streaming Multiprocessors）的概念，把四個fermi中的SM拼在一起，SMM一共有5組。同時也優化了流水線結構，減少了STALL數量等等。
有興趣的同學可以在這幾種架構白皮書中找到更多信息。
2 helloworld與軟件結構
2.1 helloworld
先看一段最基本的CUDA程序：
void CPUFunction()
{
printf("This function is defined to run on the CPU.\n");
}
__global__ void GPUFunction()
{
printf("This function is defined to run on the GPU.\n");
}
int main()
{
CPUFunction();
GPUFunction<<<1, 1>>>();
cudaDeviceSynchronize();
}
CUDA程序可以分成兩部分。一部分和C一樣，運行在CPU上；另一部分會單獨編譯，運行在“device”， 也就是GPU上。程序入口仍舊是 int main()，主要語法與C無異。
其中，__global__ void GPUFunction() 被稱作核函數（kernel function），是cuda核心程序的入口部分。程序執行至這里時，將代碼段交由GPU執行。
- 注意要加上
__global__關鍵字，表示這是運行在GPU上的函數。 - 核函數在調用時，使用三對尖括號，里面的兩個變量分別對應block數量與thread數量，引出下面的內容：
2.2 軟件結構
硬件上，運算單元被划分成了SM-->SP的層次，而相對應的在軟件上也划分了grid-->block-->thread這樣的結構：其中，block、 grid 都可以設計為三維，擁有x、y、z三個維度的信息。
軟件和硬件之間有着對應關系：
在一個核函數的運行過程當中，調度器會把Block整個地放在SM上進行運算，同時Block上的線程會運行在SM上的SP內。也就是說SM和block有對應關系，SP和thread有對應關系。但是這種映射關系並不是滿射，只能說有關聯：
首先：一個BLOCK不能分到多個SM上執行，但是不同的Block有可能會分到相同的SM（這個是調度器控制的，對用戶不可見）。既然sharedmemory在sm上，這樣就解釋了為什么Block之間沒有辦法共享shared memory，也不能夠進行同步，否則會產生死鎖。當然，根據GPU的物理結構， 還有DRAM可以共享。
block被調度至sm上后，相應的，thread與SP相對應。當然很明顯thread數量要遠遠大於SM內的Sp數量（G80內為8個），所以多個線程並不是真正的全部並行執行的，而是依靠類似操作系統中的作業調度系統，在時間片上輪轉。同一時間內，在SM上運行的thread數量就是sp數量。
grid的概念是相對於核函數來說的。由一個單獨的kernel啟動的所有線程組成一個grid，grid中所有線程共享global memory。kernel function之間也可以並行執行甚至遞歸執行，在后面會介紹。
上面的helloworld程度段只包含了一個核函數，這個核函數內只有一個block，每個block內1個thread。當然可以這樣設計：
someKernel<<<10, 1>>()配置為在 10 個線程塊（每個均具有單線程）中運行后，將運行 10 次。someKernel<<<1, 10>>()配置為在具有 10 線程的單個線程塊中運行后，將運行 10 次。someKernel<<<10, 10>>()配置為在 10 個線程塊（每個均具有 10 線程）中運行后，將運行 100 次。
2.3 SIMT
現在回頭看下前面的hello world程序段。如果把核函數的入口改成GPUFunction<<<10, 10>>>();，那么就會打印100次GPU的輸出。只編寫一行代碼，但這100個線程全部運行一段相同的程序段，這種就是CUDA的編程模型：SIMT（Single Instruction Multiple Thread）
SIMT和SIMD的概念有些容易混淆。SIMD（multiple data）指的是單指令多數據，把一個運算指令交由多個運算部件，強調使用運算部件和向量操作，來批量的對數據做處理和提升數據層面的並行性。
比如上圖的第三列。一條加法指令同時分配給4個ALU（或者是FPU， whatever）。
而SIMT強調的是線程級別，使用一條指令，運行在不同線程上，來處理不同的數據集。個人理解兩者並不是非此即彼的關系，關注的重點不太一樣。
2.3 使用自己的threadid、blockid
之前說CUDA的編程思想是SIMT，即使用一段相同的程序段（指令），為每個線程分配不同的任務。這樣就需要類似循環一樣，每一個線程都需要一個id來標識自己的唯一性。
這些變量在CUDA之中已經預先賦值，用戶可以直接調用，分別是：threadIdx blockIdx blockDim GridDim四個。其中，blockDim表示每個block中有多少個thread， gridDim表示grid之中有多少個block；剩下兩個變量標志了某一線程所處的線程號、塊號。
每一個grid、block、thread都有自己的索引（從0開始）通過將此變量與 blockIdx.x 和 threadIdx.x 變量結合使用，並借助慣用表達式 threadIdx.x + blockIdx.x * blockDim.x 在包含多個線程的多個線程塊之間組織並行執行，並行性將得以提升。
eg：
執行配置 <<<10, 10>>> 將啟動共計擁有 100 個線程的網格，這些線程均包含在由 10 個線程組成的 10 個線程塊中。因此，我們希望每個線程（0 至 99 之間）都能計算該線程的某個唯一索引。
- 如果線程塊
blockIdx.x等於0，則blockIdx.x * blockDim.x為0。向0添加可能的threadIdx.x值（0至9），之后便可在包含 100 個線程的網格內生成索引0至9。 - 如果線程塊
blockIdx.x等於1，則blockIdx.x * blockDim.x為10。向10添加可能的threadIdx.x值（0至9），之后便可在包含 100 個線程的網格內生成索引10至19。 - 如果線程塊
blockIdx.x等於5，則blockIdx.x * blockDim.x為50。向50添加可能的threadIdx.x值（0至9），之后便可在包含 100 個線程的網格內生成索引50至59。 - 如果線程塊
blockIdx.x等於9，則blockIdx.x * blockDim.x為90。向90添加可能的threadIdx.x值（0至9），之后便可在包含 100 個線程的網格內生成索引90至99。‘
當然，任務數量不可能永遠和線程數量正好匹配。（很多時候任務數量更大），這時：
-
首先需要保證線程數量>=任務數。然后可以使用if判斷條件，只給部分線程分配任務。
-
情景是，線程數往往比對應的數據數量要少，這時候就必須給每個線程分配多余1個任務。為了使得每個線程的任務盡享公平，這里采用這種跨網格的算法來實現。
具體到代碼，是這么實現的：
__global void kernel(int *a, int N)
{
int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;
int gridStride = gridDim.x * blockDim.x;
//index很好理解。就是線程號加上block號（像二維數組
//此處的步長，其實就是所有的線程數量（？）例如上圖的步長應該為8；blockdim為block內線程數量（4），griddim為grid內block數量（2）
for (int i = indexWithinTheGrid; i < N; i += gridStride)
{
// do work on a[i];
}
}
即使待處理任務數量不是總線程數的整數倍也可以完備的執行。
2.4高維度的網格、塊
可以將網格和線程塊定義為最多具有 3 個維度。使用多個維度定義網格和線程塊絕不會對其性能造成任何影響，但這在處理具有多個維度的數據時可能非常有用，例如 2D 矩陣。如要定義二維或三維網格或線程塊，可以使用 CUDA 的 dim3 類型，即如下所示：
dim3 threads_per_block(16, 16, 1);
dim3 number_of_blocks(16, 16, 1);
someKernel<<<number_of_blocks, threads_per_block>>>();
鑒於以上示例，someKernel 內部的變量 gridDim.x、gridDim.y、blockDim.x 和 blockDim.y 均將等於 16。
2.5 查詢設備信息
跟warp、deviceid等等有關的信息，被放在一個結構體prop之中（類似於句柄）。使用如下方法可以進行訪問：
int deviceId;
cudaGetDevice(&deviceId);
cudaDeviceProp props;
cudaGetDeviceProperties(&props, deviceId);
int computeCapabilityMajor = props.major;
int computeCapabilityMinor = props.minor;
int multiProcessorCount = props.multiProcessorCount;
int warpSize = props.warpSize;
printf("Device ID: %d\nNumber of SMs: %d\nCompute Capability Major: %d\nCompute Capability Minor: %d\nWarp Size: %d\n", deviceId, multiProcessorCount, computeCapabilityMajor, computeCapabilityMinor, warpSize);
3 編譯
CUDA在編譯時使用編譯器nvcc，作為一個C的擴展，nvcc的編譯方法和gcc/g++類似。
在編譯時，nvcc會用g++對其中cpu部分的代碼編譯，余下gpu的部分使用cudacc進行編譯，首先生成一個虛擬環境下的.ptx文件，之后再根據具體GPU類型生成不同的二進制碼。
使用nvcc -o指令可以直接獲得可執行文件。
4 同步問題
上面的代碼中有一行還沒有提及，即cudaDeviceSynchronize() ，用來控制CUDA上設備進行同步。
“同步”時，所有的對象會再執行到這條語句的時候進入阻塞狀態，等到其他對象都運行到了這里，再一同運行下面的指令（“對象”可以是線程、block等等）。同步在CUDA中是一個很關鍵的概念，因為如果不使用同步會導致很多意外錯誤，仍舊以上面的helloworld為例子：
void helloCPU()
{
printf("Hello from the CPU.\n");
}
__global__ void helloGPU()
{
printf("Hello also from the CPU.\n");
}
int main()
{
helloCPU();
helloGPU<<<1,1>>>();
//cudaDeviceSynchronize();
}
如果把最后的cudaDeviceSynchronize注釋掉，來自GPU的信息就不會被打印。總這里可以看出進行同步的作用：
- cudaDeviceSynchronize()：
在CPU端調用，阻塞CPU端線程的執行，一直到GPU完成之前CUDA的任務，包括kernel函數、數據拷貝等。
- _syncthreads()：block內的線程同步。
因為所有線程並不是完全並行的，有的在片上，有的在片下，因為不能保證全部線程都運行速度一樣。例如，一個很簡單的規約操作： n個數兩兩相加，最后求總和。這個過程中就需要注意線程的同步問題。
_syncthreads()和cudaThreadSynchronize() 是有區別的。_syncthreads()控制的是塊內線程，在cuda程序段內調用，而cudaThreadSynchronize() 在cpu內調用。
5 線程的調度與線程束
5.1 線程束
GPU雖然標榜自己是多線程並行，但用戶聲明的幾千個線程可能並不會完全並行化，因為一個block上的線程數量通常遠大於SM內的SP數量。
這就說到了CUDA的另一個名詞，即線程束（warps)。這個名詞可以直接望文生義：就是一把線程，一個小的線程集合，通常為32個。GPU在線程調度的時候，會由調度器（global block scheduler）每次選擇一個線程束分配到SM上。例如一個核函數定義了128個線程，那么就會划分為4個warps。只有被選擇到的線程有資格訪問運算單元，其余的線程都將處於掛起狀態，等待訪存延遲或者在就緒隊列中等待。
Warp是很重要的概念，因為一個warp是一個GPU內的基本執行單元。在代碼運行的時候，同一個warp內的32個線程都會執行相同的指令，這對訪存、分支跳轉語句等指令來說都是很關鍵的，這一部分后面再展開。
線程束的划分是固定的。假設一個block上定義了128個線程，0-31號線程一定會分配到warp1， 32-63被分配到warp2，以此類推。如果用戶聲明的線程數量並不是32的整數倍，如一個block內有100個線程，那么就會有32-（100 mod 32） 個線程處於未激活狀態，仍然會被分配到warp上。
但線程束在調度的時候沒有太大規律，其調度不是按照順序的，只要有處於Ready Queue的warps,且當前的運算單元沒被占用時，warp調度器就會調用Ready Queue的warps去執行指令。直到warps都掛起。所以，warp1 並不是一定再warps 2 后執行。在編寫程序的時候不要依賴於線程束調度順序來編寫程序，記得使用_syncthreads()來控制線程的同步，不要自行判斷線程執行的先后順序。
5.2 控制流語句
先寫下CPU中怎么處理分支跳轉語句。
為了加快CPU的運行速度，人們設計了5級的靜態流水線，分為取指(IF)、譯碼(ID)、執行(EXE)、訪存（MEM）、寫回(WB)五個階段。IF階段，CPU從指令寄存器IR中讀出下一條指令的首地址，並交給譯碼器將操作數、指令編碼進行翻譯。各操作數都准備就緒后根據操作碼送入運算單元執行，運算完成后如果是一個訪存操作需要按照地址訪問存儲器，最后再把運算結果或者訪存結果寫回寄存器。
絕大部分指令下，這個流水線是可以不間斷的執行的（遇到RAW、WAW等相關性問題需要通過重命名寄存器等方法消除，此處不再展開）。一個例外是跳轉語句。以一條misp指令BENZ R1, # 100為例，如果r1寄存器的值不為0，那么程序就會以100為偏移量進行跳轉。當流水線遇到這條指令時，需要譯碼結束，操作數都准備完成的EXE階段，才能得知這條語句是否需要跳轉以及跳轉到哪里，可問題是這時候下一條語句已經完成取指了，到底取的是哪一條指令？
早期的架構中，CPU會使用一條空轉指令NOP來填充延遲槽，遇到這條指令CPU不做任何操作，空轉一個時鍾周期等待跳轉判斷完成。可以使用其他指令填充延遲槽，但這些都需要指令調度。當然也可以使用分支預測技術。簡單來講，就是通過維持一個跳轉轉移表，來對當前語句的跳轉行為進行預測。不論最后的跳轉結果，先預判一條指令到流水線上，如果預測對了就順序執行，如果預測失敗就先把流水線全部 flush 清空，再取新的指令。
GPU的運行邏輯中也常常會遇到分支跳轉指令，然而並沒有足夠資源為每一個SP都進行分支預測，所以GPU的處理十分簡單粗暴：同一warp內所有的線程，都執行同一條指令：
1 if( threadId.x < 4) {
2 a = 0
}
3 else{
4 b= 0;
}
1 - 4 這幾行代碼，在CPU中只會執行一部分，但是GPU中會順序的執行。對於Warp1，共 0到3 號線程在邏輯上滿足if條件，4 到31號線程滿足else條件。但是warp是作為一個整體進行取指操作的，整個warp都會從1-4順序執行，只不過在運行到第二行、第四行的時候只有部分線程會執行，剩下的處於停用狀態。
可想而知，當if內代碼段很長的時候，線程空轉帶來的開銷也會越來越大。而這些都是CUDA為了保證一個線程束內的代碼執行速度嚴格一致導致的。這種問題被稱作Warp-divergence，消除 Warp-divergence 是早期GPU優化中一種很重要的優化手段。同樣需要指出的是，在有些更先進的GPU架構之中，這種由分支語句引起的性能損失也被優化掉了，這時一個Warp內線程的執行速度也不再是嚴格一致的。這一點今后如果有時間再寫一篇進行探討。
5.3 線程的選擇和切換
下面寫一些流水線上的細節，因為這一部分不同架構的區別較大，優先寫一些通用的部分。
CPU上擁有5級流水線，早期的流水線是順序執行的，也叫靜態流水。但是后來人們發現指令的執行不一定要按照順序，可以亂序發射，甚至亂序執行，再順序提交（寫回原寄存器）即可。為了實現亂序，這種動態流水線在5級流水的基礎上有了一些變化：
在取指和譯碼完成后，指令會首先進行寄存器重命名來消除之間存在的 WAW、WAR偽相關（例如： add r0, r0, r1; add r0, r2, r3 兩條指令間存在寫后寫相關。理論上第二條需要第一條寫回后才能寫回。但是通過寄存器重命名，可以先把指令1 的值存在 r5， 這樣兩條指令可以亂序執行，只要最后再按順序提交即可 ）；之后，指令被送至發射隊列中（dispatch），由發射站（dispatch port）選擇合適的指令發射出去；發射后的指令讀寄存器、執行、寫回、最后提交。
盡管目前為止，GPU上還沒有亂序的流水線，但是GPU仍有發射等等概念。指令和warp是綁定的，warp的執行要經過如下步驟：
首先， warp scheduler 會花一個時鍾周期從指令緩存中取出下一條指令（load from L1 cache ， store into cache buffer）
warp scheduler很好理解，作用在於管理warp。每次warp切換時，warp scheduler會從位於就緒隊列的warp中選擇一個發射。會有一個計分板（score boarding）來實現這些功能，決定選擇哪個warp，類似於操作系統的作業調度。
導致warp處於stall 狀態的因素會有很多，其中最常見的就是Memory Dependency：之前說過，global memory距離很遠，每次取數都要花費很長時間。訪問一次global memory 大概需要等待200個時鍾周期，在這段時間整個warp都必須進入stall狀態；除此之外warp也有可能因為別的原因阻塞，例如用戶主動控制的_syncthreads()指令；或者cache miss 導致的延遲（Instruction Fetch、Constant），等等。
指令到達發射站后，根據自身內容選擇合適的發射端口（dispatch port）。例如如果是整形數字會正常進入ALU；如果是cos、sin等三角函數，則會被發射到SFU，浮點數則進入FPU，等等。
從上述過程可以看到，指令的發射存在不少瓶頸：首先遇到的瓶頸是warp stall，之后是發射站：指令在發射站上等待，有的復雜函數指令可能還需要發射多次；然后功能部件也會阻塞，G80為例，一個SM上只有2塊SFU，如果設計大量的三角函數就只能等待。
CUDA上很有意思的一點是，warp在片上的切換時間基本可以忽略不計。CPU中，進行進程切換需要先保存上下文、處理好寄存器和程序執行信息。但GPU上這一過程幾乎不需要等待。主要的原因是線程的寄存器都是私有的（關於存儲器下一次再寫），只要分配到了某個線程，其生命周期與對應線程相同。所以位於片上的warp移至STALL狀態后並不需要做中斷和恢復。
5.4 多Block 下的 warp
討論一個特殊情況幫助理解warp的調度：
假設一個核函數定義為：someKernel<<<n, 256>>()，擁有n個block，每個block下有256個線程。在CUDA調度時，正好有3塊block被分配到了同一塊SM上，意味着這塊SM上需要運行3*256 = 768 個線程。某一時刻內，一塊SM上有幾個warp呢？
Sm下共有 768/32 = 24 個warps，但是任意時刻只會執行24中的一個warp。也就是說不同block下的warp如果被分配到了同一塊SM，也會相互等待，不能並行工作
5.5 不同GPU架構上的線程調度
隨着GPU的發展，有的GPU提供了更多的core和運算能力，不再是8個sp的結構。這種情況下前面的數字會有所不同，調度方法也會有區別。
Tesla 的 線程切換是一來SM上的MT issue部件實現的（每個warp都有自己的指令）， 同時在G80中，還有一塊4KiB 的Register File (RF)用來存放線程狀態。
Fermi下更加清晰一些：每個SM上設計了2個warp schedulers，且每個warp scheduler 對應一個發射站，每次都可以選擇一個warp，完成2條指令；在 core內有dispatch port，且浮點運算器放在了core里面。warp scheduler一個周期選擇一個warp，dispatch unit一個周期可以發射一條指令，因此只要不在其它地方阻塞，一個SM內每個周期都可以處理兩個warp。
fermi已經設計了32個core，正好和一個warp數量內thread相同。但是卻又設計了兩個warp並行執行，稱 dual scheduler，平均下來每個時鍾都可以處理一個warp。個人猜測這樣做也是為了提升並行效率，nvidia自己寫：“Using this elegant model of dual-issue, Fermi achieves near peak hardware performance.”， hhhh。
Maxwell提出了SMM概念，把四塊從前的SM拼在了一起。但是四塊小的SM仍然在資源上是有分區的，warp會屬於其中某個分區，分區之間的部分功能部件不進行共享。每一個小SM上只有一個warp scheduler，但還是保留了兩個發射單元， 這樣一來如果是普通的ALU運算指令，每個時鍾周期的warp就可以填滿整個128個core，額外的一個發射單元可以減少發射時的阻塞。
（可以翻回去看上文的圖）
6 並發CUDA流
在 CUDA 編程中，流是由按順序執行的一系列命令構成。在 CUDA 應用程序中，核函數的執行以及一些內存傳輸均在 CUDA 流中進行。一直到目前為止，還沒有直接與 CUDA 流打交道；但實際上您的 CUDA 代碼已在名為默認流的流中執行了其核函數。
這種模式下，每個核函數必須在上一個完成后才能開始。
當然也可以使用多個流。除默認流以外，CUDA 程序員還可創建並使用非默認 CUDA 流，此舉可支持執行多個操作，例如在不同的流中並發執行多個核函數。多流的使用可以為加速應用程序帶來另外一個層次的並行，並能提供更多應用程序的優化機會。
注意：
- 給定流中的所有操作會按序執行。
- 就不同非默認流中的操作而言，無法保證其會按彼此之間的任何特定順序執行。
- 默認流具有阻斷能力，即，它會等待其它已在運行的所有流完成當前操作之后才運行，但在其自身運行完畢之前亦會阻礙其它流的運行。
6.1 創建和銷毀流
注意這核函數是四個參數。而不是兩個
要在非默認CUDA流中啟動CUDA核函數，必須將流作為執行配置的第4個可選參數傳遞給該核函數。
#include <unistd.h>
cudaStream_t stream; // CUDA流的類型為 `cudaStream_t`
cudaStreamCreate(&stream); // 注意，必須將一個指針傳遞給 `cudaCreateStream`
someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>(); // `stream` 作為第4個EC參數傳遞
cudaStreamDestroy(stream); // 注意，將值（而不是指針）傳遞給 `cudaDestroyStream`
第三個參數的含義目前為止並不涉及。此參數允許程序員提供共享內存中為每個內核啟動動態分配的字節數，每個塊分配給共享內存的默認字節數為“0”。
6.2 實現並發流
例如上面的for例子，可以寫成：
#include <stdio.h>
#include <unistd.h>
__global__ void printNumber(int number)
{
printf("%d\n", number);
}
int main()
{
for (int i = 0; i < 5; ++i)
{
cudaStream_t stream;
cudaStreamCreate(&stream);
printNumber<<<1, 1, 0, stream>>>(i);
cudaStreamDestroy(stream);
}
cudaDeviceSynchronize();
}
可以看到，現在流已經變成了並行執行。