CUDA 7流簡化并發(fā)
異構(gòu)計算是指有效使用系統(tǒng)中的所有處理器，包括CPU和GPU。為此，應用程序必須在多個處理器上同時執(zhí)行功能。CUDA應用程序通過在流（按順序執(zhí)行的命令序列）中，執(zhí)行異步命令來管理并發(fā)性。不同的流可能同時執(zhí)行，或彼此相對執(zhí)行命令。
在不指定流的情況下執(zhí)行異步CUDA命令時，運行時runtime將使用默認流。在CUDA 7之前，默認流是特殊流，它與設備上的所有其他流隱式同步。
CUDA 7引入了大量強大的新功能，其中包括為每個主機線程使用獨立默認流的新選項，從而避免了對傳統(tǒng)默認流的序列化。本文將展示這如何簡化CUDA程序中內(nèi)核與數(shù)據(jù)副本之間的并發(fā)。
Asynchronous Commands in CUDA
As described by the CUDA C Programming Guide, asynchronous commands return control to the calling host thread before the device has finished the requested task (they are non-blocking). These commands are:
? Kernel launches;
? Memory copies between two addresses to the same device memory;
? Memory copies from host to device of a memory block of 64 KB or less;
? Memory copies performed by functions with the Async suffix;
? Memory set function calls.
Specifying a stream for a kernel launch or host-device memory copy is optional; you can invoke CUDA commands without specifying a stream (or by setting the stream parameter to zero). The following two lines of code both launch a kernel on the default stream.
如CUDA C編程指南所述，異步命令在設備完成請求的任務之前，將控制權(quán)返回給調(diào)用主線程（它們是非阻塞的）。這些命令是：
? 內(nèi)核啟動；
? 在兩個地址之間復制內(nèi)存到相同的設備內(nèi)存；
? 從主機到設備的內(nèi)存副本，大小為64 KB或更少；
? 帶Async后綴的功能執(zhí)行的內(nèi)存副本；
? 內(nèi)存設置函數(shù)調(diào)用。
為內(nèi)核啟動或主機設備內(nèi)存副本指定流是可選的；可以在不指定流的情況下（或通過將stream參數(shù)設置為零）調(diào)用CUDA命令。以下兩行代碼都在默認流上啟動內(nèi)核。

kernel<<< blocks, threads, bytes >>>(); // default stream
kernel<<< blocks, threads, bytes, 0 >>>(); // stream 0
The Default Stream
如果并發(fā)性對性能不重要，則默認流很有用。在CUDA 7之前，每個設備都有一個用于所有主機線程的默認流，這會導致隱式同步。如《 CUDA C編程指南》中的“隱式同步”部分所述，如果主機線程向它們之間的默認流發(fā)出任何CUDA命令，則來自不同流的兩個命令不能同時運行。
CUDA 7引入了一個新選項，即每線程默認流，它具有兩個作用。首先，它為每個主機線程提供自己的默認流。這意味著由不同的主機線程發(fā)布到默認流的命令可以同時運行。其次，這些默認流是常規(guī)流。這意味著默認流中的命令可以與非默認流中的命令同時運行。
為了使每個線程默認CUDA流7或更高版本，可以用編譯nvcc命令行選項–default-stream per-thread，或#define在CUDA_API_PER_THREAD_DEFAULT_STREAM預處理宏包括CUDA頭（前cuda.h或cuda_runtime.h）。重要的是要注意：#define CUDA_API_PER_THREAD_DEFAULT_STREAM編譯代碼時，不能使用.cu文件中的此功能，nvcc因為nvcc隱式包括cuda_runtime.h在翻譯單元的頂部。
To enable per-thread default streams in CUDA 7 and later, you can either compile with the nvcc command-line option --default-stream per-thread, or #define the CUDA_API_PER_THREAD_DEFAULT_STREAM preprocessor macro before including CUDA headers (cuda.h or cuda_runtime.h). It is important to note: you cannot use #define CUDA_API_PER_THREAD_DEFAULT_STREAM to enable this behavior in a .cu file when the code is compiled by nvcc because nvcc implicitly includes cuda_runtime.h at the top of the translation unit.
A Multi-Stream Example
Let’s look at a trivial example. The following code simply launches eight copies of a simple kernel on eight streams. We launch only a single thread block for each grid so there are plenty of resources to run multiple of them concurrently. As an example of how the legacy default stream causes serialization, we add dummy kernel launches on the default stream that do no work. Here’s the code.
看一個簡單的例子。以下代碼僅在八個流上啟動一個簡單內(nèi)核的八個副本。為每個網(wǎng)格僅啟動一個線程塊，有大量資源可以同時運行多個線程。作為傳統(tǒng)默認流如何導致序列化的一個示例，在默認流上添加了無效的虛擬內(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;

}
首先，讓我們通過不帶任何選項的編譯來檢查傳統(tǒng)行為。
nvcc ./stream_test.cu -o stream_legacy
可以在NVIDIA Visual Profiler（nvvp）中運行該程序，以獲得顯示所有流和內(nèi)核啟動的時間線。圖1顯示了在配備NVIDIA GeForce GT 750M（開普勒GPU）的Macbook Pro上生成的內(nèi)核時間軸。可以在默認流上看到虛擬內(nèi)核的極小條，以及如何導致所有其它流序列化。

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

圖2：使用新的每線程默認流選項的多流示例，該選項支持完全并發(fā)執(zhí)行。
A Multi-threading Example
Let’s look at another example, designed to demonstrate how the new default stream behavior makes it easier to achieve execution concurrency in multi-threaded applications. The following example creates eight POSIX threads, and each thread calls our kernel on the default stream and then synchronizes the default stream. (We need the synchronization in this example to make sure the profiler gets the kernel start and end timestamps before the program exits.)
另一個示例旨在演示新的默認流，如何使在多線程應用程序中更容易實現(xiàn)執(zhí)行并發(fā)。下面的示例創(chuàng)建八個POSIX線程，每個線程在默認流上調(diào)用內(nèi)核，然后同步默認流。（需要進行同步，確保事件控制器在程序退出之前，獲得內(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;

}
不進行任何編譯來測試舊版默認流的行為。
nvcc ./pthread_test.cu -o pthreads_legacy
當在nvvp中運行此命令時，看到單個流，即默認流，其中所有內(nèi)核啟動均已序列化，如圖3所示。

圖3：具有舊式默認流行為的多線程示例：所有八個線程都已序列化。
使用新的每線程默認流選項進行編譯。
nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread
圖4顯示了使用每個線程的默認流時，每個線程會自動創(chuàng)建一個新流，并且它們不會同步，因此所有八個線程的內(nèi)核同時運行。

圖4：具有每個線程默認流的多線程示例：來自所有八個線程的內(nèi)核同時運行。
Tips提示
進行并發(fā)編程時，還需要記住一些其它事項。
? 切記：對于每個線程的默認流，就同步和并發(fā)而言，每個線程中的默認流的行為與常規(guī)流相同。對于舊式默認流，情況并非如此。
? 該–default-stream選項適用于每個編譯單元，確保將其應用于nvcc需要它的所有命令行。
? cudaDeviceSynchronize()繼續(xù)使用新的每線程默認流選項同步設備上的所有內(nèi)容。如果只想同步單個流，使用cudaStreamSynchronize(cudaStream_t stream)，如第二個示例中所示。
? 從CUDA 7開始，還可以使用句柄顯式訪問每個線程的默認流cudaStreamPerThread，并且可以使用句柄訪問舊式默認流cudaStreamLegacy。注意，cudaStreamLegacy碰巧將它們混合在程序中，則仍與每個線程默認流進行隱式同步。
? 可以通過將cudaStreamNonBlocking標志傳遞到cudaStreamCreate()來創(chuàng)建不與舊式默認流同步的非阻塞流。

總結(jié)

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

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