一．異步並發執行
CUDA將以下操作公開為可以彼此並發操作的獨立任務：
主機計算；
設備計算；
從主機到設備的內存傳輸；
從設備到主機的存儲器傳輸；
在給定設備的存儲器內的存儲器傳輸；
設備之間的內存傳輸。
這些操作之間實現的並發級別將取決於設備的功能集和計算能力，如下所述。
二. 主機和設備之間的並發執行
在設備完成請求的任務之前，通過異步庫函數將控制權返回給主機線程，可以促進並發主機執行。使用異步調用，當適當的設備資源可用時，許多設備操作可以一起排隊，由CUDA驅動程序執行。這減輕了主機線程管理設備的大部分責任，讓它可以自由地執行其他任務。以下設備操作相對於主機是異步的：
內核啟動；
單個設備內存中的內存副本；
64 KB或更小的內存塊從主機到設備的內存拷貝；
由以Async為后綴的函數執行的內存拷貝；
內存設置函數調用。
程序員可以通過將CUDA啟動阻塞環境變量設置為1，全局禁用系統上運行的所有CUDA應用程序的內核啟動異步性。此功能僅用於調試目的，不應用作使生產軟件可靠運行的方法。 如果通過探查器（Nsight，Visual profiler）收集硬件計數器，則內核啟動是同步的，除非啟用了並發內核評測。如果異步內存副本涉及未頁鎖定的主機內存，則它們也將是同步的。
三. 並發內核執行
一些計算能力為2.x及更高的設備可以同時執行多個內核。應用程序可以通過檢查concurrentKernels設備屬性（請參閱設備枚舉）來查詢此功能，對於支持此功能的設備，該屬性等於1。
設備可以並發執行的最大內核啟動次數取決於其計算能力，如表15所示。
一個CUDA上下文中的內核不能與另一個CUDA上下文中的內核同時執行。
使用許多紋理或大量本地內存的內核不太可能與其他內核同時執行。
四. 數據傳輸與內核執行的重疊
有些設備可以在內核執行的同時執行到GPU或從GPU執行的異步內存復制。應用程序可以通過檢查asyncEngineCount設備屬性（請參閱設備枚舉）來查詢此功能，對於支持此功能的設備，該屬性大於零。如果復制涉及主機內存，則必須將其頁鎖定。
還可以在內核執行（在支持concurrentKernels設備屬性的設備上）和/或與設備之間的副本（對於支持asyncEngineCount屬性的設備）同時執行設備內復制。使用標准內存復制功能啟動設備內復制，目標和源地址位於同一設備上。
五. 並發數據傳輸
一些計算能力為2.x及更高的設備可以在設備之間重疊拷貝。應用程序可以通過檢查asyncEngineCount設備屬性（請參閱設備枚舉）來查詢此功能，對於支持此功能的設備，該屬性等於2。為了重疊，傳輸中涉及的任何主機內存都必須被頁鎖定。
六. 線程流
應用程序通過流管理上述並發操作。流是按順序執行的命令序列（可能由不同的主機線程發出）。另一方面，不同的流可能會執行它們的命令，彼此不按順序執行，或者同時執行；這種行為沒有保證，因此不應依賴於正確性（例如，內核間通信未定義）。流上發出的命令可以在滿足命令的所有依賴項時執行。依賴項可以是以前在同一流上啟動的命令，也可以是來自其他流的依賴項。同步調用的成功完成保證所有啟動的命令都已完成。
七. 創造與銷毀
流是通過創建一個流對象並將其指定為內核啟動序列和主機<->設備內存副本的流參數來定義的。下面的代碼示例創建兩個流並在頁鎖定內存中分配float的數組hostPtr。
cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);
以下代碼示例將這些流中的每個流定義為一個從主機到設備的內存副本、一個內核啟動和一個從設備到主機的內存副本的序列：
for (int i = 0; i < 2; ++i)
{
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}
每個流將其輸入數組hostPtr的部分復制到設備內存中的數組inputDevPtr，通過調用MyKernel（）在設備上處理inputDevPtr，並將結果outputDevPtr復制回hostPtr的相同部分。重疊行為描述了在本例中，根據設備的能力，流如何重疊。請注意，hostPtr必須指向頁鎖定的主機內存，否則將發生任何重疊。
通過調用cudaStreamDestroy（）釋放流。
for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);
如果在調用cudaStreamDestroy（）時設備仍在流中工作，則函數將立即返回，並且在設備完成流中的所有工作后，與流關聯的資源將自動釋放。
八． 默認流
內核啟動和主機<->未指定任何流參數或等效於將流參數設置為零的設備內存副本將被發送到默認流。因此，它們是按順序執行的。
對於使用--default stream per thread編譯標志編譯的代碼（或在包含CUDA頭（CUDA.h和CUDA_runtime.h）之前定義CUDA API_per_thread_default_stream宏的代碼），默認流是常規流，每個主機線程都有自己的默認流。
注意：#define CUDA_API_PER_THREAD_DEFAULT_STREAM 1不能用於在nvcc編譯代碼時啟用此行為，因為nvcc在轉換單元的頂部隱式包含CUDA_runtime.h。在這種情況下，需要使用--default stream per thread編譯標志，或者需要使用
-DCUDA_API_per_thread_default_stream=1編譯器標志定義
CUDA_API_per_thread_default_stream宏。
對於使用--default stream legacy compilation標志編譯的代碼，默認流是一個稱為空流的特殊流，每個設備都有一個用於所有主機線程的空流。空流是特殊的，因為它導致隱式同步，如隱式同步中所述。
對於未指定--default stream compilation標志而編譯的代碼，假定-default stream legacy為默認值。
九．顯式同步
有多種方法可以顯式地同步流。
cudaDeviceSynchronize（）等待，直到所有主機線程的所有流中的所有前面的命令都完成。
cudaStreamSynchronize（）接受流作為參數，並等待給定流中所有前面的命令完成。它可用於將主機與特定流同步，從而允許其他流在設備上繼續執行。 cudaStreamWaitEvent（）接受一個流和一個事件作為參數（有關事件的描述，請參見事件），並使調用cudaStreamWaitEvent（）后添加到給定流的所有命令延遲執行，直到給定事件完成。
cudaStreamQuery（）為應用程序提供了一種方法，可以知道流中所有前面的命令是否都已完成。
十．隱式同步
如果主機線程在不同流之間發出以下任一操作，則來自不同流的兩個命令不能同時運行：
頁鎖定的主機內存分配，
設備內存分配，
一個設備內存設置，
兩個地址之間到同一設備存儲器的存儲器副本，
任何對空流的CUDA命令，
在計算能力3.x和計算能力7.x中描述的L1/共享內存配置之間的切換。
對於支持並發內核執行且具有計算能力3.0或更低版本的設備，任何需要依賴性檢查以查看流式內核啟動是否完成的操作：
只有當從CUDA上下文中的任何流啟動的所有先前內核的所有線程塊都已開始執行時，才能開始執行；
阻止以后從CUDA上下文中的任何流啟動所有內核，直到選中的內核啟動完成。
需要依賴項檢查的操作包括與要檢查的啟動和對該流上cudaStreamQuery（）的任何調用相同的流中的任何其他命令。因此，應用程序應遵循以下准則，以提高並發內核執行的潛力： 所有獨立操作應在從屬操作之前發布，任何類型的同步都應盡量延遲。
十一．重疊行為
兩個流之間的執行重疊量取決於向每個流發出命令的順序，以及設備是否支持數據傳輸和內核執行重疊（請參閱數據傳輸和內核執行重疊）、並發內核執行重疊（請參閱並發內核執行），和/或並發數據傳輸（參見並發數據傳輸）。
例如，在不支持並發數據傳輸的設備上，創建和銷毀代碼樣本的兩個流根本不重疊，因為從主機到設備的內存副本在從設備到主機的內存副本頒發給流[0]之后，會將從主機到設備的內存副本頒發給流[1]，因此，它只能在從設備到主機的內存復制完成后啟動。如果代碼按以下方式重寫（並且假設設備支持數據傳輸和內核執行的重疊）
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
MyKernel<<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
然后，向流[1]發出的從主機到設備的內存復制與向流[0]發出的內核啟動重疊。
在支持並發數據傳輸的設備上，創建和銷毀的代碼示例的兩個流確實重疊：從主機到設備的內存副本發出到流[1]與從設備到主機的內存副本發出到流[0]重疊，甚至與內核啟動發出到流[0]重疊（假設設備支持數據傳輸和內核執行重疊）。但是，對於計算能力為3.0或更低的設備，內核執行不可能重疊，因為在將設備到主機的內存副本頒發給流[0]之后，第二次內核啟動將頒發給流[1]，因此它將被阻止，直到根據隱式同步完成頒發給流[0]的第一次內核啟動。如果如上所述重寫代碼，則內核執行重疊（假設設備支持並發內核執行），因為在將設備到主機的內存副本頒發給流[0]之前，第二次內核啟動被頒發給流[1]。然而，在這種情況下，根據隱式同步，從設備到主機的內存拷貝（發給流[0]）只與發給流[1]的內核啟動的最后一個線程塊重疊，后者只能表示內核總執行時間的一小部分。
十二．主機函數（回調）
運行時提供了一種通過cudaLaunchHostFunc（）將CPU函數調用插入流中的方法。一旦在回調完成之前向流發出所有命令，則在主機上執行提供的函數。
下面的代碼示例在將主機到設備的內存副本、內核啟動和設備到主機的內存副本發送到每個流之后，將主機函數MyCallback添加到兩個流中的每個流中。此函數將在每個設備到主機的內存復制完成后開始在主機上執行。
void
CUDART_CB MyCallback(cudaStream_t stream, cudaError_t status, void *data)
{
printf("Inside callback %d\n", (size_t)data);
}
...
for (size_t i = 0; i < 2; ++i)
{
cudaMemcpyAsync(devPtrIn[i], hostPtr[i], size, cudaMemcpyHostToDevice, stream[i]);
MyKernel<<<100, 512, 0, stream[i]>>>(devPtrOut[i], devPtrIn[i], size);
cudaMemcpyAsync(hostPtr[i], devPtrOut[i], size, cudaMemcpyDeviceToHost, stream[i]);
cudaLaunchHostFunc(stream[i], MyCallback, (void*)i);
}
在宿主函數之后在流中發出的命令在函數完成之前不會開始執行。
列隊到流中的主機函數不能（直接或間接）進行CUDAAPI調用，因為如果它進行這樣的調用而導致死鎖，它可能會自己等待。
十三．流優先級
流的相對優先級可以在創建時使用cudaStreamCreateWithPriority（）指定。可以使用cudaDeviceGetStreamPriorityRange（）函數獲得允許的優先級范圍，順序為[最高優先級，最低優先級]。在運行時，高優先級流中的掛起工作優先於低優先級流中的掛起工作。
下面的代碼示例獲取當前設備允許的優先級范圍，並創建具有最高和最低可用優先級的流。
// get the range of stream priorities for this device
int priority_high, priority_low;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
// create streams with highest and lowest available priorities
cudaStream_t st_high, st_low;
cudaStreamCreateWithPriority(&st_high, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&st_low, cudaStreamNonBlocking, priority_low);