以下CUDA sample是分別用C++和CUDA實現的計算一維直方圖，并對其中使用到的CUDA函數進行了解說，code參考了《GPU高性能編程CUDA實戰》一書的第九章，各個文件內容如下：

funset.cpp:

#include "funset.hpp"
#include <random>
#include <iostream>
#include <vector>
#include <memory>
#include <string>
#include <algorithm>
#include "common.hpp"
#include <opencv2/opencv.hpp>int test_calculate_histogram()
{const int length{ 10 * 1024 * 1024 }; // 100MBstd::unique_ptr<unsigned char[]> data(new unsigned char[length]);generator_random_number<unsigned char>(data.get(), length, 0, 255);const int hist_size{ 256 };std::unique_ptr<size_t[]> hist1(new size_t[hist_size]), hist2(new size_t[hist_size]);std::for_each(hist1.get(), hist1.get() + hist_size, [](size_t& n) {n = 0; });std::for_each(hist2.get(), hist2.get() + hist_size, [](size_t& n) {n = 0; });float elapsed_time1{ 0.f }, elapsed_time2{ 0.f }; // millisecondssize_t value1{ 0 }, value2{ 0 };int ret = calculate_histogram_cpu(data.get(), length, hist1.get(), value1, &elapsed_time1);if (ret != 0) PRINT_ERROR_INFO(calculate_histogram_cpu);ret = calculate_histogram_gpu(data.get(), length, hist2.get(), value2, &elapsed_time2);if (ret != 0) PRINT_ERROR_INFO(calculate_histogram_gpu);if (value1 != value2) {fprintf(stderr, "their values are different: val1: %d, val2: %d\n", value1, value2);return -1;}for (int i = 0; i < hist_size; ++i) {if (hist1[i] != hist2[i]) {fprintf(stderr, "their values are different at: %d, val1: %d, val2: %d\n",i, hist1[i], hist2[i]);return -1;}}fprintf(stderr, "test calculate histogram: cpu run time: %f ms, gpu run time: %f ms\n", elapsed_time1, elapsed_time2);return 0;
}

calculate_histogram.cpp:

#include "funset.hpp"
#include <chrono>
#include "common.hpp"int calculate_histogram_cpu(const unsigned char* data, int length, size_t* hist, size_t& value, float* elapsed_time)
{auto start = std::chrono::steady_clock::now();for (int i = 0; i < length; ++i) {++hist[data[i]];}value = 0;for (int i = 0; i < 256; ++i) {value += hist[i];}auto end = std::chrono::steady_clock::now();auto duration = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start);*elapsed_time = duration.count() * 1.0e-6;return 0;
}

calculate_histogram.cu:

#include "funset.hpp"
#include <iostream>
#include <algorithm>
#include <memory>
#include <vector>
#include <cuda_runtime.h> // For the CUDA runtime routines (prefixed with "cuda_")
#include <device_launch_parameters.h>
#include "common.hpp"/* __global__: 函數類型限定符;在設備上運行;在主機端調用,計算能力3.2及以上可以在
設備端調用;聲明的函數的返回值必須是void類型;對此類型函數的調用是異步的,即在
設備完全完成它的運行之前就返回了;對此類型函數的調用必須指定執行配置,即用于在
設備上執行函數時的grid和block的維度,以及相關的流(即插入<<<   >>>運算符);
a kernel,表示此函數為內核函數(運行在GPU上的CUDA并行計算函數稱為kernel(內核函
數),內核函數必須通過__global__函數類型限定符定義); */
__global__ static void calculate_histogram(const unsigned char* data, int length, size_t* hist)
{/* __shared__: 變量類型限定符；使用__shared__限定符，或者與__device__限定符連用，此時聲明的變量位于block中的共享存儲器空間中，與block具有相同的生命周期，僅可通過block內的所有線程訪問；__shared__和__constant__變量默認為是靜態存儲；在__shared__前可以加extern關鍵字，但表示的是變量大小由執行參數確定；__shared__變量在聲明時不能初始化；可以將CUDA C的關鍵字__shared__添加到變量聲明中，這將使這個變量駐留在共享內存中；CUDA C編譯器對共享內存中的變量與普通變量將分別采取不同的處理方式 */// clear out the accumulation buffer called temp since we are launched with// 256 threads, it is easy to clear that memory with one write per thread__shared__  size_t temp[256]; // 共享內存緩沖區temp[threadIdx.x] = 0;/* __syncthreads: 對線程塊中的線程進行同步；CUDA架構將確保，除非線程塊中的每個線程都執行了__syncthreads()，否則沒有任何線程能執行__syncthreads()之后的指令;在同一個block中的線程通過共享存儲器(sharedmemory)交換數據，并通過柵欄同步(可以在kernel函數中需要同步的位置調用__syncthreads()函數)保證線程間能夠正確地共享數據；使用clock()函數計時，在內核函數中要測量的一段代碼的開始和結束的位置分別調用一次clock()函數，并將結果記錄下來。由于調用__syncthreads()函數后，一個block中的所有thread需要的時間是相同的，因此只需要記錄每個block執行需要的時間就行了，而不需要記錄每個thread的時間 */__syncthreads();/* gridDim: 內置變量,用于描述線程網格的維度,對于所有線程塊來說,這個變量是一個常數,用來保存線程格每一維的大小,即每個線程格中線程塊的數量.一個grid最多只有二維,為dim3類型；blockDim: 內置變量,用于說明每個block的維度與尺寸.為dim3類型,包含了block在三個維度上的尺寸信息;對于所有線程塊來說,這個變量是一個常數,保存的是線程塊中每一維的線程數量;blockIdx: 內置變量,變量中包含的值就是當前執行設備代碼的線程塊的索引;用于說明當前thread所在的block在整個grid中的位置,blockIdx.x取值范圍是[0,gridDim.x-1],blockIdx.y取值范圍是[0, gridDim.y-1].為uint3類型,包含了一個block在grid中各個維度上的索引信息;threadIdx: 內置變量,變量中包含的值就是當前執行設備代碼的線程索引;用于說明當前thread在block中的位置;如果線程是一維的可獲取threadIdx.x,如果是二維的還可獲取threadIdx.y,如果是三維的還可獲取threadIdx.z;為uint3類型,包含了一個thread在block中各個維度的索引信息 */// calculate the starting index and the offset to the next block that each thread will be processingint i = threadIdx.x + blockIdx.x * blockDim.x;int stride = blockDim.x * gridDim.x;while (i < length) {/* atomicAdd: 原子操作,底層硬件將確保當執行這些原子操作時，其它任何線程都不會讀取或寫入地址addr上的值。原子函數(atomicfunction)對位于全局或共享存儲器的一個32位或64位字執行read-modify-write的原子操作。也就是說，當多個線程同時訪問全局或共享存儲器的同一位置時，保證每個線程能夠實現對共享可寫數據的互斥操作：在一個操作完成之前，其它任何線程都無法訪問此地址。之所以將這一過程稱為原子操作，是因為每個線程的操作都不會影響到其它線程。換句話說，原子操作能夠保證對一個地址的當前操作完成之前，其它線程都不能訪問這個地址。atomicAdd(addr,y)：將生成一個原子的操作序列，這個操作序列包括讀取地址addr處的值，將y增加到這個值，以及將結果保存回地址addr。 */atomicAdd(&temp[data[i]], 1);i += stride;}// sync the data from the above writes to shared memory then add the shared memory values to the values from// the other thread blocks using global memory atomic adds same as before, since we have 256 threads,// updating the global histogram is just one write per thread!__syncthreads();// 將每個線程塊的直方圖合并為單個最終的直方圖atomicAdd(&(hist[threadIdx.x]), temp[threadIdx.x]);
}int calculate_histogram_gpu(const unsigned char* data, int length, size_t* hist, size_t& value, float* elapsed_time)
{/* cudaEvent_t: CUDA event types,結構體類型, CUDA事件,用于測量GPU在某個任務上花費的時間,CUDA中的事件本質上是一個GPU時間戳,由于CUDA事件是在GPU上實現的,因此它們不適于對同時包含設備代碼和主機代碼的混合代碼計時 */cudaEvent_t start, stop;// cudaEventCreate: 創建一個事件對象,異步啟動cudaEventCreate(&start);cudaEventCreate(&stop);// cudaEventRecord: 記錄一個事件,異步啟動,start記錄起始時間cudaEventRecord(start, 0);unsigned char* dev_buffer{ nullptr };size_t* dev_hist{ nullptr };// cudaMalloc: 在設備端分配內存cudaMalloc(&dev_buffer, length);cudaMalloc(&dev_hist, 256 * sizeof(size_t));/* cudaMemcpy: 在主機端和設備端拷貝數據,此函數第四個參數僅能是下面之一:(1). cudaMemcpyHostToHost: 拷貝數據從主機端到主機端(2). cudaMemcpyHostToDevice: 拷貝數據從主機端到設備端(3). cudaMemcpyDeviceToHost: 拷貝數據從設備端到主機端(4). cudaMemcpyDeviceToDevice: 拷貝數據從設備端到設備端(5). cudaMemcpyDefault: 從指針值自動推斷拷貝數據方向,需要支持統一虛擬尋址(CUDA6.0及以上版本)cudaMemcpy函數對于主機是同步的 */cudaMemcpy(dev_buffer, data, length, cudaMemcpyHostToDevice);/* cudaMemset: 存儲器初始化函數,在GPU內存上執行。用指定的值初始化或設置設備內存 */cudaMemset(dev_hist, 0, 256 * sizeof(size_t));// cudaDeviceProp: cuda設備屬性結構體// kernel launch - 2x the number of mps gave best timingcudaDeviceProp prop;// cudaGetDeviceProperties: 獲取GPU設備相關信息cudaGetDeviceProperties(&prop, 0);// cudaDeviceProp::multiProcessorCount: 設備上多處理器的數量int blocks = prop.multiProcessorCount;fprintf(stderr, "multiProcessorCount: %d\n", blocks);/* <<< >>>: 為CUDA引入的運算符,指定線程網格和線程塊維度等,傳遞執行參數給CUDA編譯器和運行時系統,用于說明內核函數中的線程數量,以及線程是如何組織的;尖括號中這些參數并不是傳遞給設備代碼的參數,而是告訴運行時如何啟動設備代碼,傳遞給設備代碼本身的參數是放在圓括號中傳遞的,就像標準的函數調用一樣;不同計算能力的設備對線程的總數和組織方式有不同的約束;必須先為kernel中用到的數組或變量分配好足夠的空間,再調用kernel函數,否則在GPU計算時會發生錯誤,例如越界等;使用運行時API時,需要在調用的內核函數名與參數列表直接以<<<Dg,Db,Ns,S>>>的形式設置執行配置,其中：Dg是一個dim3型變量,用于設置grid的維度和各個維度上的尺寸.設置好Dg后,grid中將有Dg.x*Dg.y個block,Dg.z必須為1;Db是一個dim3型變量,用于設置block的維度和各個維度上的尺寸.設置好Db后,每個block中將有Db.x*Db.y*Db.z個thread;Ns是一個size_t型變量,指定各塊為此調用動態分配的共享存儲器大小,這些動態分配的存儲器可供聲明為外部數組(extern __shared__)的其他任何變量使用;Ns是一個可選參數,默認值為0;S為cudaStream_t類型,用于設置與內核函數關聯的流.S是一個可選參數,默認值0. */// 當線程塊的數量為GPU中處理器數量的2倍時，將達到最優性能calculate_histogram << <blocks * 2, 256 >> >(dev_buffer, length, dev_hist);cudaMemcpy(hist, dev_hist, 256 * sizeof(size_t), cudaMemcpyDeviceToHost);value = 0;for (int i = 0; i < 256; ++i) {value += hist[i];}// cudaFree: 釋放設備上由cudaMalloc函數分配的內存cudaFree(dev_buffer);cudaFree(dev_hist);// cudaEventRecord: 記錄一個事件,異步啟動,stop記錄結束時間cudaEventRecord(stop, 0);// cudaEventSynchronize: 事件同步,等待一個事件完成,異步啟動cudaEventSynchronize(stop);// cudaEventElapseTime: 計算兩個事件之間經歷的時間,單位為毫秒,異步啟動cudaEventElapsedTime(elapsed_time, start, stop);// cudaEventDestroy: 銷毀事件對象,異步啟動cudaEventDestroy(start);cudaEventDestroy(stop);return 0;
}

執行結果如下：可見使用C++和CUDA實現的結果是完全一致的：

GitHub：https://github.com/fengbingchun/CUDA_Test

總結

以上是生活随笔為你收集整理的CUDA Samples: Calculate Histogram(atomicAdd)的全部內容，希望文章能夠幫你解決所遇到的問題。

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