作者：翟天保Steven
版權(quán)聲明：著作權(quán)歸作者所有，商業(yè)轉(zhuǎn)載請聯(lián)系作者獲得授權(quán)，非商業(yè)轉(zhuǎn)載請注明出處

實現(xiàn)原理

? ? ? ?atomicMax和 atomicMin是 CUDA 中的原子操作，用于在并行計算中安全地更新共享變量的最大值和最小值。它們確保在多線程環(huán)境中，多個線程對同一個變量的訪問不會導(dǎo)致數(shù)據(jù)競爭。使用 atomicMax可以在一個線程中比較當(dāng)前值與新值，并在新值更大時更新，而 atomicMin則是用于比較和更新最小值。這些操作對于需要從多個線程中匯總結(jié)果的應(yīng)用至關(guān)重要，能夠確保最終結(jié)果的準(zhǔn)確性。

? ? ? ?本文將通過一個實戰(zhàn)案例，進行atomic求最值的展示。

? ? ? ?（注意本文案例基于OpenCV實現(xiàn)，因為我工作圍繞各類圖像展開，這樣方便些，但是對CUDA而言，核心部分與OpenCV無關(guān)，可根據(jù)自身場景和數(shù)據(jù)結(jié)構(gòu)進行更改。）

C++測試代碼

ImageProcessing.cuh

#pragma once
#include <cuda_runtime.h>
#include <iostream>
#include <vector>
#include <opencv2/opencv.hpp>
#include <device_launch_parameters.h>using namespace cv;
using namespace std;#define TILE_WIDTH 16// 預(yù)準(zhǔn)備過程
void warmupCUDA();// 圖像最值計算-CPU
void calcMaxMin_CPU(cv::Mat input, uchar &maxV, uchar &minV);// 圖像最值計算-GPU
void calcMaxMin_GPU(cv::Mat input, uchar &maxV, uchar &minV);

ImageProcessing.cu

#include "ImageProcessing.cuh"// 預(yù)準(zhǔn)備過程
void warmupCUDA()
{float* dummy_data;cudaMalloc((void**)&dummy_data, sizeof(float));cudaFree(dummy_data);
}// 圖像最值計算-CPU
void calcMaxMin_CPU(cv::Mat input, uchar &maxV, uchar &minV)
{int row = input.rows;int col = input.cols;// 初始化最值maxV = 0;minV = 255;for (int i = 0; i < row; ++i){for (int j = 0; j < col; ++j){if (input.at<uchar>(i, j) > maxV){maxV = input.at<uchar>(i, j);}if (input.at<uchar>(i, j) < minV){minV = input.at<uchar>(i, j);}}}
}// 獲取最大最小值核函數(shù)
__global__ void getMaxMinValue_CUDA(uchar* inputImage, int width, int height, int *maxV, int *minV)
{int row = blockIdx.y * blockDim.y + threadIdx.y;int col = blockIdx.x * blockDim.x + threadIdx.x;if (row < height && col < width){atomicMax(maxV, int(inputImage[row * width + col]));atomicMin(minV, int(inputImage[row * width + col]));}
}// 圖像最值計算-GPU
void calcMaxMin_GPU(cv::Mat input, uchar &maxV, uchar &minV)
{int row = input.rows;int col = input.cols;// 定義計時器float spendtime = 0.0f;cudaEvent_t start, end;cudaEventCreate(&start);cudaEventCreate(&end);// 分配GPU內(nèi)存	uchar* d_inputImage;cudaMalloc(&d_inputImage, row * col * sizeof(uchar));// 將輸入圖像數(shù)據(jù)從主機內(nèi)存復(fù)制到GPU內(nèi)存cudaMemcpy(d_inputImage, input.data, row * col * sizeof(uchar), cudaMemcpyHostToDevice);// 計算塊和線程的大小dim3 blockSize(TILE_WIDTH, TILE_WIDTH);dim3 gridSize((col + blockSize.x - 1) / blockSize.x, (row + blockSize.y - 1) / blockSize.y);// 求最值int h_maxValue = 0;int h_minValue = 255;int *d_maxValue;int *d_minValue;cudaMalloc((void**)&d_maxValue, sizeof(int));cudaMalloc((void**)&d_minValue, sizeof(int));cudaMemcpy(d_maxValue, &h_maxValue, sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(d_minValue, &h_minValue, sizeof(int), cudaMemcpyHostToDevice);getMaxMinValue_CUDA << <gridSize, blockSize >> > (d_inputImage, col, row, d_maxValue, d_minValue);cudaMemcpy(&h_maxValue, d_maxValue, sizeof(int), cudaMemcpyDeviceToHost);cudaMemcpy(&h_minValue, d_minValue, sizeof(int), cudaMemcpyDeviceToHost);maxV = uchar(h_maxValue);minV = uchar(h_minValue);
}

main.cpp

#include "ImageProcessing.cuh"void main()
{// 預(yù)準(zhǔn)備warmupCUDA();cout << "calcMaxMin test begin." << endl;// 加載cv::Mat src = imread("test pic/test5.jpg", 0);// 調(diào)整數(shù)據(jù)區(qū)間cv::Mat src2;cv::normalize(src, src2, 20, 230, NORM_MINMAX);// CPU版本clock_t s1, e1;s1 = clock();uchar maxV1, minV1;calcMaxMin_CPU(src2, maxV1, minV1);e1 = clock();cout << "CPU time:" << double(e1 - s1) << "ms" << endl;cout << "maxV1:" << int(maxV1) << endl;cout << "minV1:" << int(minV1) << endl;// GPU版本clock_t s2, e2;s2 = clock();uchar maxV2, minV2;calcMaxMin_GPU(src2, maxV2, minV2);e2 = clock();cout << "GPU time:" << double(e2 - s2) << "ms" << endl;cout << "maxV2:" << int(maxV2) << endl;cout << "minV2:" << int(minV2) << endl;cout << "calcMaxMin test end." << endl;}

測試效果?

? ? ? ?在本文案例中，我通過歸一化函數(shù)將圖像的最值設(shè)為20和230，所以驗證功能是否正確，只需要判斷下函數(shù)執(zhí)行完輸出的最值是不是20和230即可。速度方面，CUDA也是很快的，我原以為這種簡單計算CPU會更有優(yōu)勢。

? ? ? ?該功能相對簡單，但也很常用。后續(xù)我會寫一篇關(guān)于歸一化的CUDA文章，歸一化中很重要的一部分就是確認(rèn)最值。

? ? ? ?如果函數(shù)有什么可以改進完善的地方，非常歡迎大家指出，一同進步何樂而不為呢~

? ? ? ?如果文章幫助到你了，可以點個贊讓我知道，我會很快樂~加油！