CUDA编程入门

CUDA是一个并行盘算框架.用于盘算加速.是nvidia家的产物.广泛地应用于现在的深度学习加速.  
一句话形貌就是:cuda辅助我们把运算从cpu放到gpu上做,gpu多线程同时处置运算,到达加速效果.

从一个简朴例子提及:

#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;
}

这段代码很简朴,对两个数组对应位置元素相加.数组很大,有100万个元素.
CUDA编程入门
代码运行时间在0.075s.

改写代码使之运行于gpu

gpu上能够运算的函数,在cuda中我们称之为kernel.由nvcc将其编译为可以在GPU上运行的花样.

#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;
}

nvcc编译的文件的后缀为.cu

  • cuda中界说kernel在函数前加上__global声明就可以了.
  • 在显存上分配内存使用cudaMallocManaged
  • 挪用一个函数使用<<< >>>符号.好比对add的函数的挪用使用`add<<<1, 1>>>(N, x, y);`,关于其中参数的意义,后文再做注释.
  • 需要cudaDeviceSynchronize()让cpu守候gpu上的盘算做完再执行cpu上的操作

CUDA编程入门
可以用nvprof做更详细的性能剖析.   

注重用sudo 否则可能报错.
sudo /usr/local/cuda/bin/nvprof ./add_cuda

CUDA编程入门
gpu上add用了194ms.

这里,我们注重到,跑在gpu反而比cpu更慢了.由于我们这段代码里`add<<<1, 1>>>(N, x, y);`并没有施展gpu并行运算的优势,反而由于多了一些cpu与gpu的交互使得程序变慢了.

用GPU threads加速运算

重点来了
CUDA GPUS有多组Streaming Multiprocessor(SM).每个SM可以运行多个thread block. 每一个thread block有多个thread.
如下图所示:
CUDA编程入门
注重几个要害变量:

redis++:Redis持久化 rdb & aof 工作原理及流程图 (三)

  • blockDim.x 解释晰一个thread block内含有多少个thread
  • threadIdx.x 解释晰当前thread在该thread blcok内的index
  • blockIdx.x 解释晰当前是第几个thread block

我们要做的就是把盘算分配到所有的thread上去.这些thread上并行地做运算,从而到达加速的目的.

前面我们说到在cuda内挪用一个函数(称之为kernel)的用法为<<<p1,p2>>>,好比`add<<<1, 1>>>(N, x, y);` 第一个参数的寄义即为thread block的数目,第二个参数的寄义为block内介入运算的thread数目.

现在来改写一下代码:

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

// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x; 
  int stride = blockDim.x;
  printf("index=%d,stride=%d\n",index,stride);
  for (int i = index; i < n; i+=stride)
  {
    y[i] = x[i] + y[i];
    if(index == 0)
    {
        printf("i=%d,blockIdx.x=%d,thread.x=%d\n",i,blockIdx.x,threadIdx.x);
    }
  }
}

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, 256>>>(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;
}

注重add的写法,我们把0,256,512…放到thread1盘算,把1,257,…放到thread2盘算,依次类推.挪用的时刻,add<<<1, 256>>>(N, x, y);解释我们只把盘算分配到了thread block1内的256个thread去做.
编译这个程序(注重把代码里的printf注释掉,由于要统计程序运行时间):nvcc add_block.cu -o add_cuda_blcok -I/usr/local/cuda-9.0/include/ -L/usr/local/cuda-9.0/lib64
CUDA编程入门
可以看到add的gpu时间仅仅用了2.87ms

CUDA编程入门
程序的整体运行时间为0.13s,主要是cudaMallocManaged,cudaDeviceSynchronize之类的操作耗费了比较多的时间.

再一次改写代码
这一次我们用更多的thread block.

  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<<numBlocks, blockSize>>>(N, x, y);
// Kernel function to add the elements of two arrays
__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];
      //printf("i=%d,blockIdx.x=%d\n",i,blockIdx.x);
  }
}

编译:nvcc add_grid.cu -o add_cuda_grid -I/usr/local/cuda-9.0/include/ -L/usr/local/cuda-9.0/lib64
统计性能:
CUDA编程入门
可以看出来,gpu上add所用的时间进一步缩小到1.8ms

参考:https://devblogs.nvidia.com/even-easier-introduction-cuda/

原创文章,作者:28x29新闻网,如若转载,请注明出处:https://www.28x29.com/archives/4093.html