CUDA Samples: image normalize(mean/standard deviation)

42 篇文章 27 订阅

以下CUDA sample是分别用C++和CUDA实现的通过均值和标准差对图像进行类似归一化的操作,并对其中使用到的CUDA函数进行了解说,各个文件内容如下:

关于均值和标准差的计算公式可参考: http://blog.csdn.net/fengbingchun/article/details/73323475

funset.cpp: 

 

 
  1. #include "funset.hpp"

  2. #include <random>

  3. #include <iostream>

  4. #include <vector>

  5. #include <memory>

  6. #include <string>

  7. #include <algorithm>

  8. #include "common.hpp"

  9. #include <opencv2/opencv.hpp>

  10.  
  11. int test_image_normalize()

  12. {

  13. std::string image_name{ "E:/GitCode/CUDA_Test/test_images/lena.png" };

  14. cv::Mat matSrc = cv::imread(image_name);

  15. if (!matSrc.data) {

  16. fprintf(stderr, "read image fail: %s\n", image_name.c_str());

  17. return -1;

  18. }

  19.  
  20. const int width{ 511 }, height{ 473 }, channels{ 3 };

  21. cv::resize(matSrc, matSrc, cv::Size(width, height));

  22. matSrc.convertTo(matSrc, CV_32FC3);

  23. std::vector<cv::Mat> matSplit;

  24. cv::split(matSrc, matSplit);

  25. CHECK(matSplit.size() == channels);

  26. std::unique_ptr<float[]> data(new float[matSplit[0].cols * matSplit[0].rows * channels]);

  27. size_t length{ matSplit[0].cols * matSplit[0].rows * sizeof(float) };

  28. for (int i = 0; i < channels; ++i) {

  29. memcpy(data.get() + matSplit[0].cols * matSplit[0].rows * i, matSplit[i].data, length);

  30. }

  31.  
  32. float elapsed_time1{ 0.f }, elapsed_time2{ 0.f }; // milliseconds

  33. std::unique_ptr<float[]> dst1(new float[matSplit[0].cols * matSplit[0].rows * channels]);

  34. std::unique_ptr<float[]> dst2(new float[matSplit[0].cols * matSplit[0].rows * channels]);

  35.  
  36. int ret = image_normalize_cpu(data.get(), dst1.get(), width, height, channels, &elapsed_time1);

  37. if (ret != 0) PRINT_ERROR_INFO(image_normalize_cpu);

  38.  
  39. ret = image_normalize_gpu(data.get(), dst2.get(), width, height, channels, &elapsed_time2);

  40. if (ret != 0) PRINT_ERROR_INFO(image_normalize_gpu);

  41.  
  42. int count{ 0 }, num{ width * height * channels };

  43. for (int i = 0; i < num; ++i) {

  44. if (fabs(dst1[i] - dst2[i]) > 0.01/*EPS_*/) {

  45. fprintf(stderr, "index: %d, val1: %f, val2: %f\n", i, dst1[i], dst2[i]);

  46. ++count;

  47. }

  48. if (count > 100) return -1;

  49. }

  50.  
  51. std::vector<cv::Mat> merge(channels);

  52. for (int i = 0; i < channels; ++i) {

  53. merge[i] = cv::Mat(height, width, CV_32FC1, dst2.get() + i * width * height);

  54. }

  55. cv::Mat dst3;

  56. cv::merge(merge, dst3);

  57. dst3.convertTo(dst3, CV_8UC3, 255.f);

  58. cv::imwrite("E:/GitCode/CUDA_Test/test_images/image_normalize.png", dst3);

  59. //cv::resize(matSrc, matSrc, cv::Size(width, height));

  60. //cv::imwrite("E:/GitCode/CUDA_Test/test_images/image_src.png", matSrc);

  61.  
  62. fprintf(stderr, "test image normalize: cpu run time: %f ms, gpu run time: %f ms\n", elapsed_time1, elapsed_time2);

  63.  
  64. return 0;

  65. }

image_normalize.cpp:

 

 
  1. #include "funset.hpp"

  2. #include <vector>

  3. #include <chrono>

  4. #include "common.hpp"

  5.  
  6. int image_normalize_cpu(const float* src, float* dst, int width, int height, int channels, float* elapsed_time)

  7. {

  8. auto start = std::chrono::steady_clock::now();

  9.  
  10. const int offset{ width * height };

  11. for (int i = 0; i < channels; ++i) {

  12. const float* p1 = src + offset * i;

  13. float* p2 = dst + offset * i;

  14. float mean{ 0.f }, sd{ 0.f };

  15.  
  16. for (int t = 0; t < offset; ++t) {

  17. mean += p1[t];

  18. sd += pow(p1[t], 2.f);

  19. p2[t] = p1[t];

  20. }

  21.  
  22. mean /= offset;

  23. sd /= offset;

  24. sd -= pow(mean, 2.f);

  25. sd = sqrt(sd);

  26. if (sd < EPS_) sd = 1.f;

  27.  
  28. for (int t = 0; t < offset; ++t) {

  29. p2[t] = (p1[t] - mean) / sd;

  30. }

  31. }

  32.  
  33. auto end = std::chrono::steady_clock::now();

  34. auto duration = std::chrono::duration_cast<std::chrono::nanoseconds>(end - start);

  35. *elapsed_time = duration.count() * 1.0e-6;

  36.  
  37. return 0;

  38. }

image_normalize.cu:

 

 
  1. #include "funset.hpp"

  2. #include <iostream>

  3. #include <memory>

  4. #include <algorithm>

  5. #include <cmath>

  6. #include <cuda_runtime.h> // For the CUDA runtime routines (prefixed with "cuda_")

  7. #include <device_launch_parameters.h>

  8. #include "common.hpp"

  9.  
  10. /* __global__: 函数类型限定符;在设备上运行;在主机端调用,计算能力3.2及以上可以在

  11. 设备端调用;声明的函数的返回值必须是void类型;对此类型函数的调用是异步的,即在

  12. 设备完全完成它的运行之前就返回了;对此类型函数的调用必须指定执行配置,即用于在

  13. 设备上执行函数时的grid和block的维度,以及相关的流(即插入<<< >>>运算符);

  14. a kernel,表示此函数为内核函数(运行在GPU上的CUDA并行计算函数称为kernel(内核函

  15. 数),内核函数必须通过__global__函数类型限定符定义);*/

  16. __global__ static void image_normalize(const float* src, float* dst, int count, int offset)

  17. {

  18. /* gridDim: 内置变量,用于描述线程网格的维度,对于所有线程块来说,这个

  19. 变量是一个常数,用来保存线程格每一维的大小,即每个线程格中线程块的数量.

  20. 一个grid最多只有二维,为dim3类型;

  21. blockDim: 内置变量,用于说明每个block的维度与尺寸.为dim3类型,包含

  22. 了block在三个维度上的尺寸信息;对于所有线程块来说,这个变量是一个常数,

  23. 保存的是线程块中每一维的线程数量;

  24. blockIdx: 内置变量,变量中包含的值就是当前执行设备代码的线程块的索引;用

  25. 于说明当前thread所在的block在整个grid中的位置,blockIdx.x取值范围是

  26. [0,gridDim.x-1],blockIdx.y取值范围是[0, gridDim.y-1].为uint3类型,

  27. 包含了一个block在grid中各个维度上的索引信息;

  28. threadIdx: 内置变量,变量中包含的值就是当前执行设备代码的线程索引;用于

  29. 说明当前thread在block中的位置;如果线程是一维的可获取threadIdx.x,如果

  30. 是二维的还可获取threadIdx.y,如果是三维的还可获取threadIdx.z;为uint3类

  31. 型,包含了一个thread在block中各个维度的索引信息 */

  32. int index = threadIdx.x + blockIdx.x * blockDim.x;

  33. if (index > count - 1) return;

  34.  
  35. const float* input = src + index * offset;

  36. float* output = dst + index * offset;

  37. float mean{ 0.f }, sd{ 0.f };

  38.  
  39. for (size_t i = 0; i < offset; ++i) {

  40. mean += input[i];

  41. sd += pow(input[i], 2.f);

  42. output[i] = input[i];

  43. }

  44.  
  45. mean /= offset;

  46. sd /= offset;

  47. sd -= pow(mean, 2.f);

  48. sd = sqrt(sd);

  49. if (sd < EPS_) sd = 1.f;

  50.  
  51. for (size_t i = 0; i < offset; ++i) {

  52. output[i] = (input[i] - mean) / sd;

  53. }

  54. }

  55.  
  56. int image_normalize_gpu(const float* src, float* dst, int width, int height, int channels, float* elapsed_time)

  57. {

  58. /* cudaEvent_t: CUDA event types,结构体类型, CUDA事件,用于测量GPU在某

  59. 个任务上花费的时间,CUDA中的事件本质上是一个GPU时间戳,由于CUDA事件是在

  60. GPU上实现的,因此它们不适于对同时包含设备代码和主机代码的混合代码计时 */

  61. cudaEvent_t start, stop;

  62. // cudaEventCreate: 创建一个事件对象,异步启动

  63. cudaEventCreate(&start);

  64. cudaEventCreate(&stop);

  65. // cudaEventRecord: 记录一个事件,异步启动,start记录起始时间

  66. cudaEventRecord(start, 0);

  67.  
  68. float *dev_src{ nullptr }, *dev_dst{ nullptr };

  69. size_t length{ width * height * channels * sizeof(float) };

  70.  
  71. // cudaMalloc: 在设备端分配内存

  72. cudaMalloc(&dev_src, length);

  73. cudaMalloc(&dev_dst, length);

  74.  
  75. /* cudaMemcpy: 在主机端和设备端拷贝数据,此函数第四个参数仅能是下面之一:

  76. (1). cudaMemcpyHostToHost: 拷贝数据从主机端到主机端

  77. (2). cudaMemcpyHostToDevice: 拷贝数据从主机端到设备端

  78. (3). cudaMemcpyDeviceToHost: 拷贝数据从设备端到主机端

  79. (4). cudaMemcpyDeviceToDevice: 拷贝数据从设备端到设备端

  80. (5). cudaMemcpyDefault: 从指针值自动推断拷贝数据方向,需要支持

  81. 统一虚拟寻址(CUDA6.0及以上版本)

  82. cudaMemcpy函数对于主机是同步的 */

  83. cudaMemcpy(dev_src, src, length, cudaMemcpyHostToDevice);

  84.  
  85. /* <<< >>>: 为CUDA引入的运算符,指定线程网格和线程块维度等,传递执行参

  86. 数给CUDA编译器和运行时系统,用于说明内核函数中的线程数量,以及线程是如何

  87. 组织的;尖括号中这些参数并不是传递给设备代码的参数,而是告诉运行时如何

  88. 启动设备代码,传递给设备代码本身的参数是放在圆括号中传递的,就像标准的函

  89. 数调用一样;不同计算能力的设备对线程的总数和组织方式有不同的约束;必须

  90. 先为kernel中用到的数组或变量分配好足够的空间,再调用kernel函数,否则在

  91. GPU计算时会发生错误,例如越界等;

  92. 使用运行时API时,需要在调用的内核函数名与参数列表直接以<<<Dg,Db,Ns,S>>>

  93. 的形式设置执行配置,其中:Dg是一个dim3型变量,用于设置grid的维度和各个

  94. 维度上的尺寸.设置好Dg后,grid中将有Dg.x*Dg.y个block,Dg.z必须为1;Db是

  95. 一个dim3型变量,用于设置block的维度和各个维度上的尺寸.设置好Db后,每个

  96. block中将有Db.x*Db.y*Db.z个thread;Ns是一个size_t型变量,指定各块为此调

  97. 用动态分配的共享存储器大小,这些动态分配的存储器可供声明为外部数组

  98. (extern __shared__)的其他任何变量使用;Ns是一个可选参数,默认值为0;S为

  99. cudaStream_t类型,用于设置与内核函数关联的流.S是一个可选参数,默认值0. */

  100. image_normalize << < channels, 512 >> >(dev_src, dev_dst, channels, width*height);

  101.  
  102. cudaMemcpy(dst, dev_dst, length, cudaMemcpyDeviceToHost);

  103.  
  104. // cudaFree: 释放设备上由cudaMalloc函数分配的内存

  105. cudaFree(dev_src);

  106. cudaFree(dev_dst);

  107.  
  108. // cudaEventRecord: 记录一个事件,异步启动,stop记录结束时间

  109. cudaEventRecord(stop, 0);

  110. // cudaEventSynchronize: 事件同步,等待一个事件完成,异步启动

  111. cudaEventSynchronize(stop);

  112. // cudaEventElapseTime: 计算两个事件之间经历的时间,单位为毫秒,异步启动

  113. cudaEventElapsedTime(elapsed_time, start, stop);

  114. // cudaEventDestroy: 销毁事件对象,异步启动

  115. cudaEventDestroy(start);

  116. cudaEventDestroy(stop);

  117.  
  118. return 0;

  119. }

原图如下:

 

 

结果图如下:

 

 

执行结果如下:由运行时间可知,GPU要远慢于CPU,后面再对GPU实现进行优化。还有一个问题是,在Release下,执行正常,在Debug下,GPU,在调用核函数时,会在for循环中无故退出,具体原因还未知,后面待进一步分析。

 

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

  • 0
    点赞
  • 4
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值