CUDA 7 Stream流簡化并發(fā)性
異構(gòu)計算是指高效地使用系統(tǒng)中的所有處理器，包括 CPU 和 GPU 。為此，應(yīng)用程序必須在多個處理器上并發(fā)執(zhí)行函數(shù)。 CUDA 應(yīng)用程序通過在 streams 中執(zhí)行異步命令來管理并發(fā)性，這些命令是按順序執(zhí)行的。不同的流可以并發(fā)地執(zhí)行它們的命令，也可以彼此無序地執(zhí)行它們的命令。
在不指定流的情況下執(zhí)行異步 CUDA 命令時，運行時使用默認(rèn)流。在 CUDA 7 之前，默認(rèn)流是一個特殊流，它隱式地與設(shè)備上的所有其他流同步。
CUDA 7 引入了大量強大的新功能 ，包括一個新的選項，可以為每個主機線程使用獨立的默認(rèn)流，這避免了傳統(tǒng)默認(rèn)流的序列化。本文將展示如何在 CUDA 程序中簡化實現(xiàn)內(nèi)核和數(shù)據(jù)副本之間的并發(fā)。

CUDA 中的異步命令
如 CUDA C 編程指南所述，異步命令在設(shè)備完成請求的任務(wù)之前將控制權(quán)返回給調(diào)用主機線程（它們是非阻塞的）。這些命令是：
? 內(nèi)核啟動；
? 存儲器在兩個地址之間復(fù)制到同一設(shè)備存儲器；
? 從主機到設(shè)備的 64kb 或更少內(nèi)存塊的內(nèi)存拷貝；
? 由后綴為 Async 的函數(shù)執(zhí)行的內(nèi)存復(fù)制；
? 內(nèi)存設(shè)置函數(shù)調(diào)用。
為內(nèi)核啟動或主機設(shè)備內(nèi)存復(fù)制指定流是可選的；可以調(diào)用 CUDA 命令而不指定流（或通過將 stream 參數(shù)設(shè)置為零）。下面兩行代碼都在默認(rèn)流上啟動內(nèi)核。
kernel<<< blocks, threads, bytes >>>(); // default stream
kernel<<< blocks, threads, bytes, 0 >>>(); // stream 0
默認(rèn)流
在并發(fā)性對性能不重要的情況下，默認(rèn)流很有用。在 CUDA 7 之前，每個設(shè)備都有一個用于所有主機線程的默認(rèn)流，這會導(dǎo)致隱式同步。正如 CUDA C 編程指南中的“隱式同步”一節(jié)所述，如果主機線程向它們之間的默認(rèn)流發(fā)出任何 CUDA 命令，來自不同流的兩個命令就不能并發(fā)運行。
CUDA 7 引入了一個新選項， 每線程默認(rèn)流 ，它有兩個效果。首先，它為每個主機線程提供自己的默認(rèn)流。這意味著不同主機線程向默認(rèn)流發(fā)出的命令可以并發(fā)運行。其次，這些默認(rèn)流是常規(guī)流。這意味著默認(rèn)流中的命令可以與非默認(rèn)流中的命令同時運行。
要在 nvcc 7 及更高版本中啟用每線程默認(rèn)流，可以在包含 CUDA 頭（ cuda.h 或 cuda_runtime.h ）之前，使用 nvcc 命令行選項 CUDA 或 #define 編譯 CUDA_API_PER_THREAD_DEFAULT_STREAM 預(yù)處理器宏。需要注意的是：當(dāng)代碼由 nvcc 編譯時，不能使用 #define CUDA_API_PER_THREAD_DEFAULT_STREAM 在. cu 文件中啟用此行為，因為 nvcc 在翻譯單元的頂部隱式包含了 cuda_runtime.h 。
Multistream多流示例
看一個小例子。下面的代碼簡單地在八個流上啟動一個簡單內(nèi)核的八個副本。只為每個網(wǎng)格啟動一個線程塊，這樣就有足夠的資源同時運行多個線程塊。作為遺留默認(rèn)流如何導(dǎo)致序列化的示例，在默認(rèn)流上添加不起作用的虛擬內(nèi)核啟動。這是密碼。
const int N = 1 << 20;

global void kernel(float *x, int n)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
x[i] = sqrt(pow(3.14159,i));
}
}

int main()
{
const int num_streams = 8;

cudaStream_t streams[num_streams];
float *data[num_streams];for (int i = 0; i < num_streams; i++) {cudaStreamCreate(&streams[i]);cudaMalloc(&data[i], N * sizeof(float));// launch one worker kernel per streamkernel<<<1, 64, 0, streams[i]>>>(data[i], N);// launch a dummy kernel on the default streamkernel<<<1, 1>>>(0, 0);
}cudaDeviceReset();return 0;

}
首先讓檢查遺留行為，通過不帶選項的編譯。
nvcc ./stream_test.cu -o stream_legacy可以在 NVIDIA visualprofiler （ nvvp ）中運行該程序，以獲得顯示所有流和內(nèi)核啟動的時間軸。圖 1 顯示了 Macbook Pro 上生成的內(nèi)核時間線，該 Macbook Pro 帶有 NVIDIA GeForce GT 750M （一臺開普勒 GPU ）。可以看到默認(rèn)流上虛擬內(nèi)核的非常小的條，以及它們?nèi)绾螌?dǎo)致所有其他流序列化。

一個簡單的多流示例在將任何交錯內(nèi)核發(fā)送到默認(rèn)流時不會實現(xiàn)并發(fā)現(xiàn)在嘗試新的每線程默認(rèn)流。
nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread
圖 2 顯示了來自 nvvp 的結(jié)果。在這里可以看到九個流之間的完全并發(fā)：默認(rèn)流（在本例中映射到流 14 ）和創(chuàng)建的其他八個流。請注意，虛擬內(nèi)核運行得如此之快，以至于很難看到在這個圖像中默認(rèn)流上有八個調(diào)用。

圖 2 ：使用新的每線程默認(rèn)流選項的多流示例，它支持完全并發(fā)執(zhí)行。
MultiThread多線程示例
看另一個例子，該示例旨在演示新的默認(rèn)流行為如何使多線程應(yīng)用程序更容易實現(xiàn)執(zhí)行并發(fā)。下面的例子創(chuàng)建了八個 POSIX 線程，每個線程在默認(rèn)流上調(diào)用的內(nèi)核，然后同步默認(rèn)流。（需要在本例中進行同步，以確保探查器在程序退出之前獲得內(nèi)核開始和結(jié)束時間戳。）
#include <pthread.h>
#include <stdio.h>

const int N = 1 << 20;

global void kernel(float *x, int n)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
for (int i = tid; i < n; i += blockDim.x * gridDim.x) {
x[i] = sqrt(pow(3.14159,i));
}
}

void *launch_kernel(void *dummy)
{
float *data;
cudaMalloc(&data, N * sizeof(float));

kernel<<<1, 64>>>(data, N);cudaStreamSynchronize(0);return NULL;

}

int main()
{
const int num_threads = 8;

pthread_t threads[num_threads];for (int i = 0; i < num_threads; i++) {if (pthread_create(&threads[i], NULL, launch_kernel, 0)) {fprintf(stderr, "Error creating threadn");return 1;}
}for (int i = 0; i < num_threads; i++) {if(pthread_join(threads[i], NULL)) {fprintf(stderr, "Error joining threadn");return 2;}
}cudaDeviceReset();return 0;

}
首先，編譯時不使用任何選項來測試遺留的默認(rèn)流行為。
nvcc ./pthread_test.cu -o pthreads_legacy
在 nvvp 中運行它時，看到一個流，默認(rèn)流，所有內(nèi)核啟動都序列化，如圖 3 所示。

圖 3 ：一個具有遺留默認(rèn)流行為的多線程示例：所有八個線程都被序列化。
讓用新的 per-thread default stream 選項編譯它。
nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread
圖 4 顯示，對于每個線程的默認(rèn)流，每個線程都會自動創(chuàng)建一個新的流，它們不會同步，因此所有八個線程的內(nèi)核都會并發(fā)運行。

圖 4 ：每個線程默認(rèn)流的多線程示例：所有八個線程的內(nèi)核同時運行。
更多提示
在為并發(fā)進行編程時，還需要記住以下幾點。
? 記住：對于每線程的默認(rèn)流，每個線程中的默認(rèn)流的行為與常規(guī)流相同，只要同步和并發(fā)就可以了。對于傳統(tǒng)的默認(rèn)流，這是不正確的。
? --default-stream 選項是按編譯單元應(yīng)用的，確保將其應(yīng)用于所有需要它的 nvcc 命令行。
? cudaDeviceSynchronize() 繼續(xù)同步設(shè)備上的所有內(nèi)容，甚至使用新的每線程默認(rèn)流選項。如果只想同步單個流，請使用 cudaStreamSynchronize(cudaStream_t stream) ，如的第二個示例所示。
? 從 CUDA 7 開始，還可以使用句柄 cudaStreamPerThread 顯式地訪問每線程的默認(rèn)流，也可以使用句柄 cudaStreamLegacy 訪問舊的默認(rèn)流。請注意， cudaStreamLegacy 仍然隱式地與每個線程的默認(rèn)流同步，如果碰巧在一個程序中混合使用它們。
? 可以通過將 cudaStreamCreate() 標(biāo)志傳遞給 cudaStreamCreate() 來創(chuàng)建不與傳統(tǒng)默認(rèn)流同步的 非阻塞流 。
立即下載 CUDA 7rc CUDA toolkitversion7 . 0 的發(fā)布候選者今天可以向 NVIDIA 注冊的開發(fā)人員提供。如果不是注冊開發(fā)者， 在 NVIDIA 開發(fā)區(qū)注冊免費訪問 。了解 這里是 CUDA 7 的特點。
想了解更多關(guān)于 Tesla 平臺上的加速計算以及使用 CUDA 進行 GPU 計算的信息嗎？參加 3 月 17 日至 20 日在圣何塞會議中心舉行的 GPU 技術(shù)會議 ，這是世界上規(guī)模最大、最重要的 GPU 開發(fā)者大會。 Parallel Forall 的讀者可以使用折扣代碼 GM15PFAB 獲得任何會議通行證 20% 的折扣！`

總結(jié)

以上是生活随笔為你收集整理的CUDA 7 Stream流简化并发性的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。