文章目錄

• 前言
• 1. 向量點乘
• 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：紋理內存

---

前言

這里繼續跟大家分享CUDA編程的知識，這次就不介紹新知識了，主要是應用之前學的CUDA編程知識實現一個小例子：向量點乘運算

---

1. 向量點乘

這里有兩個向量，$v_1=[a_1,a_2,...,a_n]$和$v_2=[b_1,b_2,...,b_n]$,向量點乘的計算公式如下：
$result={\sum }_{i=1}^n(a_i?b_i)$

---

2. C++ CUDA實現向量點乘

這里使用共享內存實現向量點乘，主要是因為共享內存的訪問速度比全局內存快一個數量級，使用多塊多線程來加速計算過程，在源碼中通過注釋來解讀實現過程。

詳細代碼如下所示：

#include <iostream> #include <stdio.h> #include <cuda.h> #include <cuda_runtime.h> #include <device_launch_parameters.h>#define N 1024 #define THREADS_PER_BLOCK 512__global__ void gpu_dot(float* d_a, float* d_b, float* d_c) {//為每個塊定義共享內存，不同塊的共享內存沒法通信__shared__ float partial_sum[THREADS_PER_BLOCK];//整個GPU設備的線程索引int tid = threadIdx.x + blockIdx.x * blockDim.x;//當前塊的線程索引int index = threadIdx.x;//每個線程進行一個元素乘法操作，//比如v1=[a1,a2,a3], v2=[b1,b2,b3]//線程t1: a1*b1, 線程t2：a2*b2, ...float sum = 0;while (tid < N){//sum += d_a[tid] * d_b[tid];//多塊多線程運行時，這里會造成計算結果出錯sum = d_a[tid] * d_b[tid];tid += blockDim.x + gridDim.x;}//將上述部分和保存到共享內存中partial_sum[index] = sum;//線程同步，等待塊中所有線程都執行完上述代碼行__syncthreads();//使用當前塊的共享內存中保存的數據循環計算部分和，//（1）將共享內存數組中數據二分，后面的數據累加到前面的數據中，//（2）然后舍棄后面的數據，在前面的數據中繼續（1）的過程int i = blockDim.x / 2;while (i != 0){if (index < i)//每次僅在二分線程中的前半部分線程數進行累加{partial_sum[index] += partial_sum[index + i];}//等待所有線程執行完上述過程__syncthreads();i /= 2;}//將當前塊的部分和計算結果保存到全局內存中if (index == 0){d_c[blockIdx.x] = partial_sum[0];} }int main() {float* h_a, *h_b, h_c, *partial_sum;float *d_a, *d_b, *d_partial_sum;//計算塊的數量和線程數量int block_calc = (N + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;int blocksPerGrid = block_calc > 32 ? 32 : block_calc;//在CPU上分配內存h_a = (float*)malloc(N * sizeof(float));h_b = (float*)malloc(N * sizeof(float));partial_sum = (float*)malloc(blocksPerGrid * sizeof(float));//在GPU上分配內存cudaMalloc(&d_a, N * sizeof(float));cudaMalloc(&d_b, N * sizeof(float));cudaMalloc(&d_partial_sum, blocksPerGrid * sizeof(float));//填充數據for (int i = 0; i < N; i++){h_a[i] = 1;h_b[i] = 2;}//將數據從CPU拷貝到GPU上cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);//調用內核函數，模塊數為blocksPerGrid, 每個模塊的線程數為THREADS_PER_BLOCKgpu_dot << <blocksPerGrid, THREADS_PER_BLOCK >> > (d_a, d_b, d_partial_sum);//將結果從GPU上拷貝到CPU上cudaMemcpy(partial_sum, d_partial_sum, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost);//計算最終的向量點乘結果：累加每個塊的計算結果h_c = 0;for (int i = 0; i < blocksPerGrid; i++){h_c += partial_sum[i];}printf("GPU向量點乘計算結果為：%f\n", h_c);//釋放GPU內存cudaFree(d_a);cudaFree(d_b);cudaFree(d_partial_sum);//釋放GPU內存free(h_a);free(h_b);free(partial_sum);system("pause");return 0; }

---

3. 執行結果

上述例子中兩個輸入向量一個為全1值，另一個為全2值，向量長度為1024，所以執行結果如下所示：

---

總結

這里并沒有給出GPU上加速效果，大家如果感興趣可以自行實現在CPU上執行的速度，因本人還是菜鳥一枚，上述過程如有錯誤，還請大神指正，謝謝！

學習資料

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

總結

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

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