本文介绍 2019_12_13_11_CUDA编程实战:初入江湖

2019_12_13_11_CUDA编程实战:初入江湖

本文由在当地较为英俊的男子金天大神原创,版权所有,欢迎转载,本文首发地址 https://jinfagang.github.io 。但请保留这段版权信息,多谢合作,有任何疑问欢迎通过微信联系我交流:jintianiloveu

CUDA看来是必备的技能了。因为在很多地方会需要,比如编写pytorch自定义层,编写TensorRT的plugin的时候,都会用到cuda编程。但其实这门技术入门比较难。本篇文章将从一个很微小的角度,来传授大家这门手艺。

看完这篇教程之后,你应该可以学会如何自己编写kernel来写一个nms的操作了。

凡事需要有一个正确的领导,同时需要遵循循序渐进的不变法则,我们列一个提纲,一个任务一个任务来完成。

  • 先学会CUDA的hello world;
  • 学会如何用cudaMalloc开辟内存;
  • 学会如何拷贝数据从cpu到gpu;
  • 学会如何拷贝数据从gpu到cpu;
  • 第一个带参数的kernel;
  • 开始逆天杀神。

GPU和CPU的数据拷贝

#include <cuda.h>
#include <cuda_runtime.h>
#include <vector>
#include <iostream>
#include <math.h>

int main(void)
{   
  
  float dets[6][4] = {
    {23, 34, 56, 76},
    {11, 23, 45, 45},
    {12, 22, 47, 47},
    {9, 45, 56, 65},
    {20, 37, 55, 75},
  };

  std::cout << sizeof(dets) << std::endl;

  float *dev_dets, *dev_scores;
  cudaMemcpy(dev_dets, dets, sizeof(dets), cudaMemcpyHostToDevice);
  std::cout << "copied data to GPU.\n";

  // get back copied cuda data
  float *host_dets;
  cudaMemcpy(host_dets, dev_dets, sizeof(dets), cudaMemcpyDeviceToHost);
  std::cout << "copied from cuda back to host.\n";
  for (int i=0;i<sizeof(dets)/sizeof(float);i++) {
    std::cout << host_dets[i] << " ";
  }

}
	

上面这个小代码的作用就是,将dets这个二维的数组,拷贝到GPU,然后我们再从GPU拷贝回CPU,并且查看数据是否正确。

当你打印的时候,你会发现,数据是拷贝过去了,但是拷贝回来的时候数值不对了:

copied data to GPU.
copied from cuda back to host.
194582 -4.09377 4.97457e-07 -1.54942e+26 -403.131 5.08783e+10 -128.165 1.21334e-38 0 -2.18535e-33 131241 -nan -6.19379e-31 1.17302e-38 0.068985 2.8026e-44 -nan 3.04014e-20 1.79366e-43 -2.85228e-39 0 8.70721e-31 2.88443e-41 3.33458e-38 %

这是什么原因造成的呢?

原因在于,我们在将数据从host拷贝到GPU上的时候,并没有事先开辟内存空间,这其实也是很多人容易犯的一个错误,在实际调用 cudaMemcpy之前,需要先开辟内存:

float dets[6][4] = {
    {23, 34, 56, 76},
    {11, 23, 45, 45},
    {12, 22, 47, 47},
    {9, 45, 56, 65},
    {20, 37, 55, 75},
  };
  float scores[6] = {0.7, 0.6, 0.8, 0.4, 0.2, 0.6};
  float iou_threshold = 0.2;
  // copy data to gpu
  std::cout << sizeof(dets) << std::endl;
  std::cout << sizeof(scores) << std::endl;

  float *dev_dets, *dev_scores;
  cudaError_t err = cudaSuccess;
  err = cudaMalloc((void **)&dev_scores, sizeof(scores));
  err = cudaMalloc((void **)&dev_dets, sizeof(dets));
  if (err != cudaSuccess) {
    printf("cudaMalloc failed!");
    return 1;
  }
  cudaMemcpy(dev_dets, dets, sizeof(dets), cudaMemcpyHostToDevice);
  cudaMemcpy(dev_scores, scores, sizeof(scores), cudaMemcpyHostToDevice);
  std::cout << "copied data to GPU.\n";

  // get back copied cuda data
  float host_dets[sizeof(dets)/sizeof(float)];
  float host_scores[6];
  cudaMemcpy(&host_dets, dev_dets, sizeof(dets), cudaMemcpyDeviceToHost);
  cudaMemcpy(&host_scores, dev_scores, sizeof(scores), cudaMemcpyDeviceToHost);
  std::cout << "copied from cuda back to host.\n";
  std::cout << "host_dets size: " << sizeof(host_dets) << std::endl;
  for (int i=0;i<sizeof(dets)/sizeof(float);i++) {
    std::cout << host_dets[i] << " ";
  }
  std::cout << std::endl;
  for (int i=0;i<sizeof(scores)/sizeof(float);i++) {
    std::cout << static_cast<float>(host_scores[i]) << " ";
  }
  std::cout << std::endl;

  cudaFree(dev_dets);
  cudaFree(dev_scores);

  std::cout << "done.\n";

将上面的代码,修改成这样,就没错了。上面的代码,我们演示了两个东西:

  • 将CPU的数据,拷贝到了GPU;
  • 将GPU的数据又拷贝回了CPU。

大家需要注意的是,如果你直接通过取值符号试图获取GPU上指针所指向的具体数值,那么你会得到一个段错误,显然你没有办法直接从GPU地址获取它的值,除非你将那一部分数据拷贝到CPU上。所以啊,大家经常用cuda的时候,你如果要把一个tensor转为numpy,你需要首先 a.cpu().numpy()其实这个操作背后,就是拷贝一份数据到CPU的操作。

尝试对数据进行操作

接下来我们来尝试一个更难的操作,操作什么呢?我们想办法对上面的scores这个数据进行排序。其实排序很简单了,但是我们不仅仅要对scores排序,我们还希望拿到排序之前对应的下标。

这里就需要引入一个叫做 thrust的库了。这个库要说名气也没啥名气,但是你可以把它看作是英伟达官方的,CUDA界的 STL。标准模板库。这里面提供很多类似于C++的模板操作,那么为什么有了STL还需要这么一个东西呢?本质原因是,我们运算的时候,所有的数据,大部分实在GPU上运算的,当然Thrust也支持CPU类型的数据运算。此时C++的标准模板库事肯定用不了的,必须要上thrust了。

上面提出的这个问题操作,解决办法有很多种,其中一种方法可以这样:

thrust::device_vector<int> sorted_indices(sizeof(scores)/sizeof(float));
thrust::sequence(sorted_indices.begin(), sorted_indices.end(), 0);
thrust::sort_by_key(thrust::device, dev_scores, dev_scores+sizeof(scores)/sizeof(float), sorted_indices.begin());
printf("sorted done.\n");
cudaMemcpy(&host_scores, dev_scores, sizeof(scores), cudaMemcpyDeviceToHost);
for (int i=0;i<sizeof(scores)/sizeof(float);i++) {
std::cout << static_cast<float>(host_scores[i]) << " ";
}
std::cout << std::endl;
for(auto index: sorted_indices) {
std::cout << index << " ";
}
std::cout << std::endl;

很多人看这个会很头痛,为什么呢?因为它实在是太底层了,全都是内存级别的操作。但是你理解了这基本的原理,它的运算过程也就很好理解了。这个的计算结果应嘎是:

sorted done.
0.2 0.4 0.6 0.6 0.7 0.8 
4 3 1 5 0 2 
done.

也就是对scores进行排序,然后返回排序之后对应的原数据的indices。

最后扔一个问题:如果我们像从大到小排序而不是默认的从小到大,应该怎么搞?