異構(gòu)計(jì)算是指高效地使用系統(tǒng)中的所有處理器，包括 CPU 和 GPU 。為此，應(yīng)用程序必須在多個(gè)處理器上并發(fā)執(zhí)行函數(shù)。 CUDA 應(yīng)用程序通過(guò)在 streams 中執(zhí)行異步命令來(lái)管理并發(fā)性，這些命令是按順序執(zhí)行的。不同的流可以并發(fā)地執(zhí)行它們的命令，也可以彼此無(wú)序地執(zhí)行它們的命令。[見帖子[See the post 如何在 CUDA C / C ++中實(shí)現(xiàn)數(shù)據(jù)傳輸?shù)闹丿B ]

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

CUDA 7 引入了大量強(qiáng)大的新功能 ，包括一個(gè)新的選項(xiàng)，可以為每個(gè)主機(jī)線程使用獨(dú)立的默認(rèn)流，這避免了傳統(tǒng)默認(rèn)流的序列化。在這篇文章中，我將向您展示如何在 CUDA 程序中簡(jiǎn)化實(shí)現(xiàn)內(nèi)核和數(shù)據(jù)副本之間的并發(fā)。

CUDA 中的異步命令

如 CUDA C 編程指南所述，異步命令在設(shè)備完成請(qǐng)求的任務(wù)之前將控制權(quán)返回給調(diào)用主機(jī)線程（它們是非阻塞的）。這些命令是：

• 內(nèi)核啟動(dòng)；
• 存儲(chǔ)器在兩個(gè)地址之間復(fù)制到同一設(shè)備存儲(chǔ)器；
• 從主機(jī)到設(shè)備的 64kb 或更少內(nèi)存塊的內(nèi)存拷貝；
• 由后綴為 Async 的函數(shù)執(zhí)行的內(nèi)存復(fù)制；
• 內(nèi)存設(shè)置函數(shù)調(diào)用。

為內(nèi)核啟動(dòng)或主機(jī)設(shè)備內(nèi)存復(fù)制指定流是可選的；您可以調(diào)用 CUDA 命令而不指定流（或通過(guò)將 stream 參數(shù)設(shè)置為零）。下面兩行代碼都在默認(rèn)流上啟動(dòng)內(nèi)核。

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

默認(rèn)流

在并發(fā)性對(duì)性能不重要的情況下，默認(rèn)流很有用。在 CUDA 7 之前，每個(gè)設(shè)備都有一個(gè)用于所有主機(jī)線程的默認(rèn)流，這會(huì)導(dǎo)致隱式同步。正如 CUDA C 編程指南中的“隱式同步”一節(jié)所述，如果主機(jī)線程向它們之間的默認(rèn)流發(fā)出任何 CUDA 命令，來(lái)自不同流的兩個(gè)命令就不能并發(fā)運(yùn)行。

CUDA 7 引入了一個(gè)新選項(xiàng)， 每線程默認(rèn)流 ，它有兩個(gè)效果。首先，它為每個(gè)主機(jī)線程提供自己的默認(rèn)流。這意味著不同主機(jī)線程向默認(rèn)流發(fā)出的命令可以并發(fā)運(yùn)行。其次，這些默認(rèn)流是常規(guī)流。這意味著默認(rèn)流中的命令可以與非默認(rèn)流中的命令同時(shí)運(yùn)行。

要在 nvcc 7 及更高版本中啟用每線程默認(rèn)流，您可以在包含 CUDA 頭（ cuda.h 或 cuda_runtime.h ）之前，使用 nvcc 命令行選項(xiàng) CUDA 或 #define 編譯 CUDA_API_PER_THREAD_DEFAULT_STREAM 預(yù)處理器宏。需要注意的是：當(dāng)代碼由 nvcc 編譯時(shí)，不能使用 #define CUDA_API_PER_THREAD_DEFAULT_STREAM 在。 cu 文件中啟用此行為，因?yàn)?nvcc 在翻譯單元的頂部隱式包含了 cuda_runtime.h 。

多流示例

讓我們看一個(gè)小例子。下面的代碼簡(jiǎn)單地在八個(gè)流上啟動(dòng)一個(gè)簡(jiǎn)單內(nèi)核的八個(gè)副本。我們只為每個(gè)網(wǎng)格啟動(dòng)一個(gè)線程塊，這樣就有足夠的資源同時(shí)運(yùn)行多個(gè)線程塊。作為遺留默認(rèn)流如何導(dǎo)致序列化的示例，我們?cè)谀J(rèn)流上添加了不起作用的虛擬內(nèi)核啟動(dòng)。這是密碼。

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;
}

首先讓我們檢查遺留行為，通過(guò)不帶選項(xiàng)的編譯。

nvcc ./stream_test.cu -o stream_legacy

我們可以在 NVIDIA visualprofiler （nvvp）中運(yùn)行該程序，以獲得顯示所有流和內(nèi)核啟動(dòng)的時(shí)間軸。圖 1 顯示了 Macbook Pro 上生成的內(nèi)核時(shí)間線，該 Macbook Pro 帶有 NVIDIA GeForce GT 750M （一臺(tái)開普勒 GPU ）。您可以看到默認(rèn)流上虛擬內(nèi)核的非常小的條，以及它們?nèi)绾螌?dǎo)致所有其他流序列化。

現(xiàn)在讓我們嘗試新的每線程默認(rèn)流。

nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread

圖 2 顯示了來(lái)自nvvp的結(jié)果。在這里您可以看到九個(gè)流之間的完全并發(fā)：默認(rèn)流（在本例中映射到流 14 ）和我們創(chuàng)建的其他八個(gè)流。請(qǐng)注意，虛擬內(nèi)核運(yùn)行得如此之快，以至于很難看到在這個(gè)圖像中默認(rèn)流上有八個(gè)調(diào)用。

多線程示例

讓我們看另一個(gè)例子，該示例旨在演示新的默認(rèn)流行為如何使多線程應(yīng)用程序更容易實(shí)現(xiàn)執(zhí)行并發(fā)。下面的例子創(chuàng)建了八個(gè) POSIX 線程，每個(gè)線程在默認(rèn)流上調(diào)用我們的內(nèi)核，然后同步默認(rèn)流。（我們需要在本例中進(jìn)行同步，以確保探查器在程序退出之前獲得內(nèi)核開始和結(jié)束時(shí)間戳。）

#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;
}

首先，讓我們編譯時(shí)不使用任何選項(xiàng)來(lái)測(cè)試遺留的默認(rèn)流行為。

nvcc ./pthread_test.cu -o pthreads_legacy

當(dāng)我們?cè)?code style="font-size:inherit;color:inherit;margin:0px;padding:0px;border:0px;font-style:inherit;font-variant:inherit;font-weight:inherit;line-height:inherit;vertical-align:baseline;background-color:rgb(244,244,244);">nvvp中運(yùn)行它時(shí)，我們看到一個(gè)流，默認(rèn)流，所有內(nèi)核啟動(dòng)都序列化，如圖 3 所示。

讓我們用新的 per-thread default stream 選項(xiàng)編譯它。

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

圖 4 顯示，對(duì)于每個(gè)線程的默認(rèn)流，每個(gè)線程都會(huì)自動(dòng)創(chuàng)建一個(gè)新的流，它們不會(huì)同步，因此所有八個(gè)線程的內(nèi)核都會(huì)并發(fā)運(yùn)行。

更多提示

在為并發(fā)進(jìn)行編程時(shí)，還需要記住以下幾點(diǎn)。

記?。簩?duì)于每線程的默認(rèn)流，每個(gè)線程中的默認(rèn)流的行為與常規(guī)流相同，只要同步和并發(fā)就可以了。對(duì)于傳統(tǒng)的默認(rèn)流，這是不正確的。

--default-stream 選項(xiàng)是按編譯單元應(yīng)用的，因此請(qǐng)確保將其應(yīng)用于所有需要它的 nvcc 命令行。

cudaDeviceSynchronize（） 繼續(xù)同步設(shè)備上的所有內(nèi)容，甚至使用新的每線程默認(rèn)流選項(xiàng)。如果您只想同步單個(gè)流，請(qǐng)使用 cudaStreamSynchronize（cudaStream_t stream） ，如我們的第二個(gè)示例所示。

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

您可以通過(guò)將 cudaStreamCreate（） 標(biāo)志傳遞給 cudaStreamCreate（） 來(lái)創(chuàng)建不與傳統(tǒng)默認(rèn)流同步的 非阻塞流 。

關(guān)于作者

Mark Harris 是 NVIDIA 杰出的工程師，致力于 RAPIDS 。 Mark 擁有超過(guò) 20 年的 GPUs 軟件開發(fā)經(jīng)驗(yàn)，從圖形和游戲到基于物理的模擬，到并行算法和高性能計(jì)算。當(dāng)他還是北卡羅來(lái)納大學(xué)的博士生時(shí)，他意識(shí)到了一種新生的趨勢(shì)，并為此創(chuàng)造了一個(gè)名字： GPGPU （圖形處理單元上的通用計(jì)算）。

審核編輯：郭婷