異構計算是指高效地使用系統中的所有處理器，包括 CPU 和 GPU 。為此，應用程序必須在多個處理器上并發執行函數。 CUDA 應用程序通過在 streams 中執行異步命令來管理并發性，這些命令是按順序執行的。不同的流可以并發地執行它們的命令，也可以彼此無序地執行它們的命令。[見帖子[See the post 如何在 CUDA C / C ++中實現數據傳輸的重疊 ]

在不指定流的情況下執行異步 CUDA 命令時，運行時使用默認流。在 CUDA 7 之前，默認流是一個特殊流，它隱式地與設備上的所有其他流同步。

CUDA 7 引入了大量強大的新功能 ，包括一個新的選項，可以為每個主機線程使用獨立的默認流，這避免了傳統默認流的序列化。在這篇文章中，我將向您展示如何在 CUDA 程序中簡化實現內核和數據副本之間的并發。

CUDA 中的異步命令

如 CUDA C 編程指南所述，異步命令在設備完成請求的任務之前將控制權返回給調用主機線程（它們是非阻塞的）。這些命令是：

• 內核啟動；
• 存儲器在兩個地址之間復制到同一設備存儲器；
• 從主機到設備的 64kb 或更少內存塊的內存拷貝；
• 由后綴為 Async 的函數執行的內存復制；
• 內存設置函數調用。

為內核啟動或主機設備內存復制指定流是可選的；您可以調用 CUDA 命令而不指定流（或通過將 stream 參數設置為零）。下面兩行代碼都在默認流上啟動內核。

  kernel<<< blocks, threads, bytes >>>();    // default stream
  kernel<<< blocks, threads, bytes, 0 >>>(); // stream 0

默認流

在并發性對性能不重要的情況下，默認流很有用。在 CUDA 7 之前，每個設備都有一個用于所有主機線程的默認流，這會導致隱式同步。正如 CUDA C 編程指南中的“隱式同步”一節所述，如果主機線程向它們之間的默認流發出任何 CUDA 命令，來自不同流的兩個命令就不能并發運行。

CUDA 7 引入了一個新選項， 每線程默認流 ，它有兩個效果。首先，它為每個主機線程提供自己的默認流。這意味著不同主機線程向默認流發出的命令可以并發運行。其次，這些默認流是常規流。這意味著默認流中的命令可以與非默認流中的命令同時運行。

要在 nvcc 7 及更高版本中啟用每線程默認流，您可以在包含 CUDA 頭（ cuda.h 或 cuda_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 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 stream
        kernel<<<1, 64, 0, streams[i]>>>(data[i], N);

        // launch a dummy kernel on the default stream
        kernel<<<1, 1>>>(0, 0);
    }

    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)
{
    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中運行它時，我們看到一個流，默認流，所有內核啟動都序列化，如圖 3 所示。

圖 3 ：一個具有遺留默認流行為的多線程示例：所有八個線程都被序列化。

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

nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread

圖 4 顯示，對于每個線程的默認流，每個線程都會自動創建一個新的流，它們不會同步，因此所有八個線程的內核都會并發運行。

圖 4 ：每個線程默認流的多線程示例：所有八個線程的內核同時運行。

更多提示

在為并發進行編程時，還需要記住以下幾點。

記住：對于每線程的默認流，每個線程中的默認流的行為與常規流相同，只要同步和并發就可以了。對于傳統的默認流，這是不正確的。

--default-stream 選項是按編譯單元應用的，因此請確保將其應用于所有需要它的 nvcc 命令行。

cudaDeviceSynchronize（） 繼續同步設備上的所有內容，甚至使用新的每線程默認流選項。如果您只想同步單個流，請使用 cudaStreamSynchronize（cudaStream_t stream） ，如我們的第二個示例所示。

從 CUDA 7 開始，您還可以使用句柄 cudaStreamPerThread 顯式地訪問每線程的默認流，也可以使用句柄 cudaStreamLegacy 訪問舊的默認流。請注意， cudaStreamLegacy 仍然隱式地與每個線程的默認流同步，如果您碰巧在一個程序中混合使用它們。

您可以通過將 cudaStreamCreate（） 標志傳遞給 cudaStreamCreate（） 來創建不與傳統默認流同步的 非阻塞流 。

關于作者

Mark Harris 是 NVIDIA 杰出的工程師，致力于 RAPIDS 。 Mark 擁有超過 20 年的 GPUs 軟件開發經驗，從圖形和游戲到基于物理的模擬，到并行算法和高性能計算。當他還是北卡羅來納大學的博士生時，他意識到了一種新生的趨勢，并為此創造了一個名字： GPGPU （圖形處理單元上的通用計算）。

審核編輯：郭婷