CUDA 7 Stream流簡化併發性

2022-02-03 13:43:31 字數 3585 閱讀 8864

cuda 7 stream流簡化併發性

異構計算是指高效地使用系統中的所有處理器,包括 cpu 和 gpu 。為此,應用程式必須在多個處理器上併發執行函式。 cuda 應用程式通過在 streams 中執行非同步命令來管理併發性,這些命令是按順序執行的。不同的流可以併發地執行它們的命令,也可以彼此無序地執行它們的命令。

在不指定流的情況下執行非同步 cuda 命令時,執行時使用預設流。在 cuda 7 之前,預設流是乙個特殊流,它隱式地與裝置上的所有其他流同步。

cuda 7 引入了大量強大的新功能 ,包括乙個新的選項,可以為每個主機執行緒使用獨立的預設流,這避免了傳統預設流的序列化。本文將展示如何在 cuda 程式中簡化實現核心和資料副本之間的併發。

如 cuda c 程式設計指南所述,非同步命令在裝置完成請求的任務之前將控制權返回給呼叫主機執行緒(它們是非阻塞的)。這些命令是:

為核心啟動或主機裝置記憶體複製指定流是可選的;可以呼叫 cuda 命令而不指定流(或通過將 stream 引數設定為零)。下面兩行**都在預設流上啟動核心。

kernel<<< blocks, threads, bytes >>>();    // default stream
kernel<<< blocks, threads, bytes, 0 >>>(); // stream 0
在併發性對效能不重要的情況下,預設流很有用。在 cuda 7 之前,每個裝置都有乙個用於所有主機執行緒的預設流,這會導致隱式同步。正如 cuda c 程式設計指南中的「隱式同步」一節所述,如果主機執行緒向它們之間的預設流發出任何 cuda 命令,來自不同流的兩個命令就不能併發執行。

cuda 7 引入了乙個新選項, 每執行緒預設流 ,它有兩個效果。首先,它為每個主機執行緒提供自己的預設流。這意味著不同主機執行緒向預設流發出的命令可以併發執行。其次,這些預設流是常規流。這意味著預設流中的命令可以與非預設流中的命令同時執行。

要在nvcc7 及更高版本中啟用每執行緒預設流,可以在包含 cuda 頭(cuda.hcuda_runtime.h)之前,使用nvcc命令列選項 cuda 或#define編譯cuda_api_per_thread_default_stream預處理器巨集。需要注意的是:當**由nvcc編譯時,不能使用#define cuda_api_per_thread_default_stream在. cu 檔案中啟用此行為,因為nvcc在翻譯單元的頂部隱式包含了cuda_runtime.h

看乙個小例子。下面的**簡單地在八個流上啟動乙個簡單核心的八個副本。只為每個網格啟動乙個執行緒塊,這樣就有足夠的資源同時執行多個執行緒塊。作為遺留預設流如何導致序列化的示例,在預設流上新增不起作用的虛擬核心啟動。這是密碼。

const int n = 1 << 20;
__global__ void kernel(float *x, int n)
}
int main()
cudadevicereset();
return 0;
}
首先讓檢查遺留行為,通過不帶選項的編譯。

nvcc ./stream_test.cu -o stream_legacy
可以在 nvidia visualprofiler (nvvp)中執行該程式,以獲得顯示所有流和核心啟動的時間軸。圖 1 顯示了 macbook pro 上生成的核心時間線,該 macbook pro 帶有 nvidia geforce gt 750m (一台克卜勒 gpu )。可以看到預設流上虛擬核心的非常小的條,以及它們如何導致所有其他流序列化。

乙個簡單的多流示例在將任何交錯核心傳送到預設流時不會實現併發

現在嘗試新的每執行緒預設流。

nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread
圖 2 顯示了來自nvvp的結果。在這裡可以看到九個流之間的完全併發:預設流(在本例中對映到流 14 )和建立的其他八個流。請注意,虛擬核心執行得如此之快,以至於很難看到在這個影象中預設流上有八個呼叫。

圖 2 :使用新的每執行緒預設流選項的多流示例,它支援完全併發執行。

看另乙個例子,該示例旨在演示新的預設流行為如何使多執行緒應用程式更容易實現執行併發。下面的例子建立了八個 posix 執行緒,每個執行緒在預設流上呼叫的核心,然後同步預設流。(需要在本例中進行同步,以確保探查器在程式退出之前獲得核心開始和結束時間戳。)

#include
#include
const int n = 1 << 20;
__global__ void kernel(float *x, int n)
}
void *launch_kernel(void *dummy)
int main()
}
for (int i = 0; i < num_threads; i++)
}
cudadevicereset();
return 0;
}
首先,編譯時不使用任何選項來測試遺留的預設流行為。

nvcc ./pthread_test.cu -o pthreads_legacy
nvvp中執行它時,看到乙個流,預設流,所有核心啟動都序列化,如圖 3 所示。

圖 3 :乙個具有遺留預設流行為的多執行緒示例:所有八個執行緒都被序列化。

讓用新的 per-thread default stream 選項編譯它。

nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread
圖 4 顯示,對於每個執行緒的預設流,每個執行緒都會自動建立乙個新的流,它們不會同步,因此所有八個執行緒的核心都會併發執行。

圖 4 :每個執行緒預設流的多執行緒示例:所有八個執行緒的核心同時執行。

在為併發進行程式設計時,還需要記住以下幾點。

cuda toolkitversion7 . 0 的發布候選者今天可以向 nvidia 註冊的開發人員提供。如果不是註冊開發者, 在 nvidia 開發區註冊免費訪問 。了解 這裡是 cuda 7 的特點。

cuda 使用多個stream

對stream的介紹,使用兩個流 include define n 1024 1024 每次從cpu傳輸到gpu的資料塊大小 define m n 20 cpu上的總資料量 測試裝置是否支援邊執行核函式邊複製資料 bool support overlap global void add int a,...

Stream流 方法引用

流式思想 類似於 生產流水線 模型就是每乙個位置 延遲方法 返回型別是stream介面自身型別的方法,可以鏈式程式設計 包括filter,map,skip等 終結方法 返回型別不再是stream介面自身型別的方法,包括count,foreach stream流常用方法 void foreach co...

Stream流的使用

本文簡單記述stream的使用 第一種 最常用的,從資料庫查出資料後,需要封裝成vo類,把需要的字段返回給前端 listorderlist orderservice.list listlist orderlist.stream map order collect collectors.tolist ...