CUDA编程简要介绍 (一)

总说

主要是 https://devblogs.nvidia.com/even-easier-introduction-cuda/ 的简要信息提取

准备工作

我们写的cuda程序需要nvcc来进行编译, 而其实nvcc就是在安装好的cuda目录的bin文件夹下:
我们首先将相应目录添加进.bashrc里面.

 # CUDA PATH
 export PATH="/usr/local/cuda-9.2/bin:$PATH"
 export LD_LIBRARY_PATH="/usr/local/cuda-9.2/lib64:$LD_LIBRARY_PATH"

source .bashrc 进行更新.
接下来进入正题.

C++程序与CUDA程序的对比

新建add.cpp

#include <iostream>
#include <math.h>

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 1M elements

  float *x = new float[N];
  float *y = new float[N];

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the CPU
  add(N, x, y);

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  delete [] x;
  delete [] y;

  return 0;
}

对应的CUDA程序是add.cu

#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y;

  // Allocate Unified Memory – accessible from CPU or GPU
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the GPU
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  cudaDeviceSynchronize();

  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  cudaFree(x);
  cudaFree(y);
  
  return 0;
}

几个知识点:

1. 我们在add这个函数上面加了__global__分类符, 表示CUDA C++编译器这个函数可以在GPU上跑, 而且是由CPU进行调用的.
2. cydaMallocManaged()函数可以申请Unified Memory, 而在这个统一内存可以被GPU和CPU都可以使用.
3. cudaFree()用来释放内存
4. cudaDeviceSynchronize()是用来同步一下, 即等待gpu计算完成. 再执行后面的代码
5. add<<<1, 1>>>(N, x, y) 这个是什么?

<<<1, 1>>>这个是什么?

简单来说就是, 要调用CUDA核, 就要用<<< >>>来进行. 那么这两个1是什么东西呢? 其实表示的是用多少个并行线程来进行处理.
第一个1表示,用一个线程block, 第二个1表示每个线程block使用一个线程.

cuda程序的编译以及性能测试

nvcc add.cu -o add_cuda
./add_cuda
Max error: 0.000000

现在这里只有用一个线程进行实验, 我们可以看到, 因为for循环中, 是对整个数组都逐个进行相加, 如果我们直接用多个线程来做, 就相当于每个线程都会对所有的元素进行操作, 就会造成 race condition的情况.

nvprof ./add_cuda
==3355== NVPROF is profiling process 3355, command: ./add_cuda
Max error: 0
==3355== Profiling application: ./add_cuda
==3355== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
100.00%  463.25ms         1  463.25ms  463.25ms  463.25ms  add(int, float*, float*)

现在我们要进行用多个线程进行操作, 需要改写一下for循环即可, 从而让每个线程处理不同的数据

// 假设一个block我们有256个线程
// 那么我们需要numBlocks个blocks
int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);


__global__
void add(int n, float *x, float *y)
{
  // 下面三句可以认为是定式
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}

没错, 这是通用的写法,
我们来看看CUDA并行分层
enmm, 一个block里面有多个threads, 那么多个block组成一个grid. 没啥毛病.

• gridDim.x: 一个grid的dim是多少, 即包含多少个block
• gridIdx.x: 第几个grid
• blockDim.x: 一个block包含多少个threads
• blockIdx.x: 这是第几个block
• threadIdx.x: 第几个thread
  所以再看看刚才的,

  // 下面三句可以认为是定式
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)

完全没毛病, 值得注意的是, 怎么理解这个?

因为我也是初学, 所以它可能是这样工作的.
首先, 因为外面调用这个kernel是 <<<nBlocks, nTheads>>>来的. 那么表示一下子会用nBlocks * nThreads这么多个线程进行工作. 那么, 对于单个线程进入这个函数时, 函数内部通过 index来获得这个线程所处理的元素是x[i], 其中i=index. 显然啊, 这个index是看这个线程所在的block属于我们申请的所有的block中的第几个, 以及这个线程是block中的第几个线程. 因此index = blockIdx.x * blockDim.x + threadIdx.x.
申请的所有的线程同时进行相同的工作, 而for循环中的index恰好可以把每个线程处理哪个元素进行很好的对应分配, 因此经过一次运算时间, 有nBlocks * nThreads个元素已经处理完了. 自然我们要转向下一组数据, 即此时有int stride = blockDim.x * gridDim.x以及i = i+stride.

从代码上看, 这个blockIdx.x很有可能是从0(申请的blocks的第一个)开始的, 不过可能也不一定, 这个以后更多理解之后, 再看看是否正确!

Time(%)      Time     Calls       Avg       Min       Max  Name
100.00%  94.015us         1  94.015us  94.015us  94.015us  add(int, float*, float*)

到现在为止, 我们就可以告一段落.

总结

简单来说, 知道了__global__分类符是用来修饰函数, 表示该函数(称为kernel)可以被GPU调用. 然后还知道了 CUDA 的for循环的编程3句定式, 还知道了CUDA申请内存以及释放的相关函数等.