文章目錄

• 前言
• 1. CUDA流的理解
• 2. C++ 實現CUDA流
• 3. 執行結果
• 總結
• 學習資料

VS2017 CUDA編程學習1：CUDA編程兩變量加法運算
VS2017 CUDA編程學習2：在GPU上執行線程
VS2017 CUDA編程學習3：CUDA獲取設備上屬性信息
VS2017 CUDA編程學習4：CUDA并行處理初探 - 向量加法實現
VS2017 CUDA編程學習5：CUDA并行執行-線程
VS2017 CUDA編程學習6: GPU存儲器架構
VS2017 CUDA編程學習7：線程同步-共享內存
VS2017 CUDA編程學習8：線程同步-原子操作
VS2017 CUDA編程學習9：常量內存
VS2017 CUDA編程學習10：紋理內存
VS2017 CUDA編程學習實例1：CUDA實現向量點乘
VS2017 CUDA編程學習11：CUDA性能測量

---

前言

今天跟大家分享CUDA編程中一個重要概念：CUDA流。

---

1. CUDA流的理解

之前介紹的GPU內核函數，通過多線程實現加速，可以理解為數據并行。而CUDA流可以理解為任務并行，即多個相互獨立的內核函數同時執行。

CUDA流是GPU上的工作隊列，隊列中的工作將以特定的順序執行。隊列中的工作主要包括：內核函數的調用， cudaMemcpy系列的數據傳輸，以及CUDA事件的操作。CUDA流中工作的順序以添加的順序來執行，先加入的先執行，即每個CUDA流中的工作是串行執行的。

---

2. C++ 實現CUDA流

這里以GPU上加法操作為例實現CUDA流。

例子中主要調用一下API:
cudaHostAlloc/cudaFreeHost(): 使用頁面鎖定內存為CPU設備變量分配內存，頁面鎖定內存，簡單理解就是內存數據永遠存儲在內存上，不會轉儲到硬盤空間來暫存，缺點是如果內存一直不釋放，將造成內存高使用率，可能影響其他程序的運行；

cudaMemcpyAsync:用于異步傳輸數據，和cudaMemcpy函數相比，多了最后一個參數 - 指定一個特定的CUDA流，從而在這個特定的CUDA流中進行存儲器傳輸操作。這里注意，異步操作是指當該函數返回時，實際的存儲器傳輸操作可能尚未完成，甚至尚未開始。

這里要注意，每個CUDA流中的工作是串行的，而流和流之間則默認不保證順序。每個流程都有傳輸和計算工作，這樣不同的流間計算和傳輸工作可以并行執行，起到加速的作用，進而提升程序性能。

cudaStreamSynchronize(): 確保CUDA流中所有操作都結束后才返回控制權給主機（CPU）。

詳細代碼如下所示：

#include <stdio.h> #include <iostream> #include <cuda.h> #include <cuda_runtime.h> #include <cuda_runtime_api.h> #include <device_launch_parameters.h>#define N 50000__global__ void gpuAdd(int* d_a, int* d_b, int* d_c) {int tid = threadIdx.x + blockIdx.x * blockDim.x;while (tid < N){d_c[tid] = d_a[tid] + d_b[tid];tid += blockDim.x * gridDim.x;} }int main() {//定義CPU變量int* h_a, *h_b, *h_c;//為CUDA流 stream0定義GPU變量int* d_a0, *d_b0, *d_c0;//為CUDA流 stream1定義GPU變量int *d_a1, *d_b1, *d_c1;//定義CUDA流變量cudaStream_t stream0, stream1;//創建CUDA流cudaStreamCreate(&stream0);cudaStreamCreate(&stream1);//定義CUDA事件cudaEvent_t e_start, e_stop;//創建CUDA事件cudaEventCreate(&e_start);cudaEventCreate(&e_stop);//開始記錄CUDA事件cudaEventRecord(e_start, 0);//分配CPU內存，使用頁面鎖定內存, 不會發生換頁操作（將暫時不用的數據從內存中轉存硬盤上），數據永遠存在內存中cudaHostAlloc(&h_a, 2 * N * sizeof(int), cudaHostAllocDefault);cudaHostAlloc(&h_b, 2 * N * sizeof(int), cudaHostAllocDefault);cudaHostAlloc(&h_c, 2 * N * sizeof(int), cudaHostAllocDefault);//分配GPU內存cudaMalloc(&d_a0, N * sizeof(int));cudaMalloc(&d_b0, N * sizeof(int));cudaMalloc(&d_c0, N * sizeof(int));cudaMalloc(&d_a1, N * sizeof(int));cudaMalloc(&d_b1, N * sizeof(int));cudaMalloc(&d_c1, N * sizeof(int));//初始化CPU數據for (int i = 0; i < 2 * N; i++){h_a[i] = 2 * i * i;h_b[i] = i;}//為兩個CUDA流執行異步內存拷貝操作cudaMemcpyAsync(d_a0, h_a, N * sizeof(int), cudaMemcpyHostToDevice, stream0);cudaMemcpyAsync(d_a1, h_a + N, N * sizeof(int), cudaMemcpyHostToDevice,stream1);cudaMemcpyAsync(d_b0, h_b, N * sizeof(int), cudaMemcpyHostToDevice,stream0);cudaMemcpyAsync(d_b1, h_b + N, N * sizeof(int), cudaMemcpyHostToDevice,stream1);//為兩個CUDA流分別調用內核函數gpuAdd << <512, 512, 0, stream0 >> > (d_a0, d_b0, d_c0);gpuAdd << <512, 512, 0, stream1 >> > (d_a1, d_b1, d_c1);//從GPU拷貝數據到CPUcudaMemcpyAsync(h_c, d_c0, N * sizeof(int), cudaMemcpyDeviceToHost,stream0);cudaMemcpyAsync(h_c + N, d_c1, N * sizeof(int), cudaMemcpyDeviceToHost,stream1);//等待GPU所有線程運行結束cudaDeviceSynchronize();//等待所有CUDA流中工作完成cudaStreamSynchronize(stream0);cudaStreamSynchronize(stream1);//結束CUDA事件記錄cudaEventRecord(e_stop, 0);//等待CUDA事件記錄完成cudaEventSynchronize(e_stop);//GPU上運行時間統計float elapsedTime;cudaEventElapsedTime(&elapsedTime, e_start, e_stop);printf("GPU上加法操作 %d 個數耗時：%3.1f ms\n", 2 * N, elapsedTime);//驗證GPU上運行結果的準確性int correct = 1;for (int i = 0; i < 2 * N; i++){if (h_a[i] + h_b[i] != h_c[i]){correct = 0;break;}}if (correct == 1){printf("GPU上加法操作結果準確！\n");}else{printf("GPU上加法操作結果不準確！\n");}//釋放GPU內存cudaFree(d_a0);cudaFree(d_a1);cudaFree(d_b0);cudaFree(d_b1);cudaFree(d_c0);cudaFree(d_c1);//釋放CPU頁面鎖定內存cudaFreeHost(h_a);cudaFreeHost(h_b);cudaFreeHost(h_c);system("pause");return 0; }

---

3. 執行結果

---

總結

CUDA流，說白了就是實現了多任務并行的功能。今天就分享到這里，本人還是小菜鳥，如果上面描述有問題，還請大神不吝指正！

學習資料

《基于GPU加速的計算機視覺編程》

總結

以上是生活随笔為你收集整理的VS2017 CUDA编程学习12：CUDA流的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。