当前位置:   article > 正文

Cuda 编程入门_csdn xiaohu2022

csdn xiaohu2022

这篇文章写的特别好,https://blog.csdn.net/xiaohu2022/article/details/79599947,我基本就是参考这篇文章

最难理解的部门就是Grid和Block的概念,具体关系参见下图

è¿éåå¾çæè¿°

GPU之所以处理图片速度快,就是因为可以把图片的每一块区域分配给一个Thread来处理,每个Thread只负责处理图像的一块区域,当每个Thread都结束的时候,整张图就处理完成了,避免了大量for循环的使用

Grid和Block实际上组成了一张大网,当这张大网在图片上“扫一遍”之后,整张图片也就处理完了。尤其是当图片尺寸比较小,GPU资源丰富时,更是一次性就可以扫完。

要获取每个Thread对应的位置,可以使用

  1. int row = threadIdx.y + blockIdx.y * blockDim.y;
  2. int col = threadIdx.x + blockIdx.x * blockDim.x;

 举个简单例子,假如我的图片大小是12*12,如果对这144个像素点,每个像素点做归一化,且一次完成,则可以设置

  1. dim3 gridSize(3, 3);
  2. dim3 blockSize(4, 4);

这样实际就有12*12个Thread,每个Thread处理一个像素值,一次就完成了。

当输入图片尺寸很大,无法一次完成时,每个Thread就需要处理多个像素点,这时就需要“stride”了,比如此时图片大小变为24*24,gridSize和blockSize的值不变,则此时每个Thread需要处理4个像素点,比如处理(0,0)点的Thread,还要处理(0,12),(12,0),(12,12)这几个点。

下面例子是我自己编写的第一个CUDA代码,输入为一张1920*1080的热力图,输出为这张图的局部极值点,即改点的值比上下左右四个点的值都大。首先定义kernel_fun,然后在host上采用kernel_fun<<< grid, block >>>(prams...)的形式调用

 

  1. // 获取矩阵A的(row, col)元素
  2. __device__ float getElement(float *A, int height, int width, int row, int col)
  3. {
  4. return A[row * width + col];
  5. }
  6. // 为矩阵A的(row, col)元素赋值
  7. __device__ void setElement(bool *A, int height, int width, int row, int col, bool value)
  8. {
  9. A[row * width + col] = value;
  10. }
  11. __global__ void getPeak(float *d_in, int height, int width, float thre, bool *d_out)
  12. {
  13. int row = threadIdx.y + blockIdx.y * blockDim.y;
  14. int col = threadIdx.x + blockIdx.x * blockDim.x;
  15. int stride_row = blockDim.y * gridDim.y;
  16. int stride_col = blockDim.x * gridDim.x;
  17. int max_iter_row = (height + stride_row - 1) / stride_row;
  18. int max_iter_col = (width + stride_col - 1) / stride_col;
  19. for(int iter_row = 0; iter_row <= max_iter_row ; iter_row++)
  20. {
  21. col = threadIdx.x + blockIdx.x * blockDim.x;
  22. for (int iter_col = 0; iter_col <= max_iter_col; iter_col++)
  23. {
  24. //setElement(d_out, height, width, row, col, 1);
  25. if (row - 1 <= 0 || col - 1 <= 0) continue;
  26. if (row + 1 >= height || col + 1 >= width) continue;
  27. float map_cen_val = getElement(d_in, height, width, row, col);
  28. if (
  29. (map_cen_val >= thre) &&
  30. (map_cen_val >= getElement(d_in, height, width, row - 1, col)) &&
  31. (map_cen_val >= getElement(d_in, height, width, row + 1, col)) &&
  32. (map_cen_val >= getElement(d_in, height, width, row, col - 1)) &&
  33. (map_cen_val >= getElement(d_in, height, width, row, col + 1))
  34. )
  35. {
  36. setElement(d_out, height, width, row, col, true);
  37. }
  38. col += stride_col;
  39. }
  40. row += stride_row;
  41. }
  42. }

Host上的调用方法为

  1. void findPeakPose(float* input_data, int height, int width, float thre, std::vector<Point>& peak_points)
  2. {
  3. int nBytes = width * height * sizeof(float);
  4. int nBytes_output = width * height * sizeof(bool);
  5. // 申请托管内存
  6. float *d_in;
  7. bool *d_out;
  8. cudaMalloc((void**)&d_in, nBytes);
  9. cudaMalloc((void**)&d_out, nBytes_output);
  10. // 将host数据拷贝到device
  11. cudaMemcpy((void*)d_in, (void*)input_data, nBytes, cudaMemcpyHostToDevice);
  12. // 定义kernel的执行配置
  13. dim3 blockSize(32, 32);
  14. dim3 gridSize(5, 5);
  15. // 执行kernel
  16. getPeak<<<gridSize, blockSize>>>(d_in, height, width, thre, d_out);
  17. // 同步device 保证结果能正确访问
  18. cudaDeviceSynchronize();
  19. bool* output_data = (bool*)malloc(nBytes_output);
  20. cudaMemcpy((void*)output_data, (void*)d_out, nBytes_output, cudaMemcpyDeviceToHost);
  21. for(int i = 0; i < width * height; i++)
  22. {
  23. if (output_data[i])
  24. {
  25. Point p;
  26. p.width = i%1920;
  27. p.height = i/1920;
  28. peak_points.push_back(p);
  29. }
  30. }
  31. cudaFree(d_in);
  32. cudaFree(d_out);
  33. free(output_data);
  34. }

简单说明一下__device__ __global__等函数的意义(引用自https://blog.csdn.net/xiaohu2022/article/details/79599947

  • __global__:在device上执行,从host中调用(一些特定的GPU也可以从device上调用),返回类型必须是void,不支持可变参数参数,不能成为类成员函数。注意用__global__定义的kernel是异步的,这意味着host不会等待kernel执行完就执行下一步。
  • __device__:在device上执行,单仅可以从device中调用,不可以和__global__同时用。
  • __host__:在host上执行,仅可以从host上调用,一般省略不写,不可以和__global__同时用,但可和__device__,此时函数会在device和host都编译。
声明:本文内容由网友自发贡献,不代表【wpsshop博客】立场,版权归原作者所有,本站不承担相应法律责任。如您发现有侵权的内容,请联系我们。转载请注明出处:https://www.wpsshop.cn/w/Cpp五条/article/detail/468255
推荐阅读
相关标签
  

闽ICP备14008679号