苏州手机网站建设服务,wordpress向微信群发送内容,官网cms系统,培训校园网站建设简报作者#xff1a;翟天保Steven 版权声明#xff1a;著作权归作者所有#xff0c;商业转载请联系作者获得授权#xff0c;非商业转载请注明出处
实现原理 atomicMax和 atomicMin是 CUDA 中的原子操作#xff0c;用于在并行计算中安全地更新共享变量的最大值和最小值。它们确…作者翟天保Steven 版权声明著作权归作者所有商业转载请联系作者获得授权非商业转载请注明出处
实现原理 atomicMax和 atomicMin是 CUDA 中的原子操作用于在并行计算中安全地更新共享变量的最大值和最小值。它们确保在多线程环境中多个线程对同一个变量的访问不会导致数据竞争。使用 atomicMax可以在一个线程中比较当前值与新值并在新值更大时更新而 atomicMin则是用于比较和更新最小值。这些操作对于需要从多个线程中汇总结果的应用至关重要能够确保最终结果的准确性。 本文将通过一个实战案例进行atomic求最值的展示。 注意本文案例基于OpenCV实现因为我工作围绕各类图像展开这样方便些但是对CUDA而言核心部分与OpenCV无关可根据自身场景和数据结构进行更改。
C测试代码
ImageProcessing.cuh
#pragma once
#include cuda_runtime.h
#include iostream
#include vector
#include opencv2/opencv.hpp
#include device_launch_parameters.husing namespace cv;
using namespace std;#define TILE_WIDTH 16// 预准备过程
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// 预准备过程
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.atuchar(i, j) maxV){maxV input.atuchar(i, j);}if (input.atuchar(i, j) minV){minV input.atuchar(i, j);}}}
}// 获取最大最小值核函数
__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内存 uchar* d_inputImage;cudaMalloc(d_inputImage, row * col * sizeof(uchar));// 将输入图像数据从主机内存复制到GPU内存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.cuhvoid main()
{// 预准备warmupCUDA();cout calcMaxMin test begin. endl;// 加载cv::Mat src imread(test pic/test5.jpg, 0);// 调整数据区间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;}
测试效果 在本文案例中我通过归一化函数将图像的最值设为20和230所以验证功能是否正确只需要判断下函数执行完输出的最值是不是20和230即可。速度方面CUDA也是很快的我原以为这种简单计算CPU会更有优势。 该功能相对简单但也很常用。后续我会写一篇关于归一化的CUDA文章归一化中很重要的一部分就是确认最值。 如果函数有什么可以改进完善的地方非常欢迎大家指出一同进步何乐而不为呢~ 如果文章帮助到你了可以点个赞让我知道我会很快乐~加油
