默認流
對于需要指定 cudaStream_t
參數的 cuda API
,如果將 0
作為實參傳入,則視為使用默認流;對于不需要指定 cudaStream_t
參數的 cuda API
,則也視為使用默認流。
在 cuda
中,默認流有兩種類型,一種是 legacy
默認流,一種是 per-thread
默認流,這兩種默認流的同步行為不一樣,在使用的時候需要注意一下。具體使用哪種默認流,有以下三種方式進行指定:
- 不指定,默認使用
legacy
默認流; - 在編譯
cuda
程序時,通過nvcc
的--default-stream
進行指定,可選的取值是:{legacy|per-thread}
; - 在需要傳入
cudaStream_t
/CUstream
參數的時候,指定:cudaStreamLegacy
/CU_STREAM_LEGACY
(legacy
默認流),或cudaStreamPerThread
/CU_STREAM_PER_THREAD
(per-thread
默認流)。
下面重點介紹一下legacy
默認流和per-thread
默認流在同步行為上的差異。
legacy 默認流
legacy
默認流會與同一個 CUcontext
(如果是使用 runtime
的 API
,則每個設備對應一個 CUcontext
)中的其他流都進行同步,但 non-blocking
的流除外。也即是,在執行 legacy
默認流中的任務之前,會先等待其他所有 blocking
的流執行完成,然后開始執行 legacy
默認流中的任務,并且在 legacy
默認流后面出現的其他 blocking
流中的任務,會先等待 legacy
默認流中的任務執行完成,再開始執行。
假設,下面的代碼在流 s
中 launch
了核函數 k_1
,在 legacy
默認流中 launch
了核函數 k_2
,在流 s
中 launch
了核函數 k_3
:
cudaStream_t s;
cudaStreamCreate(&s);
k_1<<<1, 1, 0, s>>>();
k_2<<<1, 1>>>();
k_3<<<1, 1, 0, s>>>();
則上述代碼的同步行為是 k_2
會被 k_1
阻塞,k_3
會被 k_2
阻塞。
但如果是 non-blocking
流,則不會出現上述同步行為,也即是,下面三個核函數會存在并行執行的情況,例如:
cudaStream_t s;
cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
k_1<<<1, 1, 0, s>>>();
k_2<<<1, 1>>>();
k_3<<<1, 1, 0, s>>>();
per-thread 默認流
在一個 CUcontext
內,每個線程都有一個 per-thread
默認流,這個流不與其他流進行同步(就像是一個顯式創建的流那樣),但如果在一個程序中同時使用了 legacy
默認流和 per-thread
默認流,則 per-thread
默認流會與 legacy
默認流保持同步。