CUDA加速学习

今天学习了一下NVIDIA官方教程第一课

CUDA 提供了一种可扩展 C、C++、Python 和 Fortran 等语言的编码范式

我对 CUDA理解，CUDA就像是C的超集一样，提供了对GPU的操作，我目前觉得CUDA没有太多面向对象的东西，更多的是面向过程的代码

CPU上的代码称为主机代码，而在GPU上的代码称为设备代码，GPU就可以理解成一种外设一样，专门用来处理一些计算

在调用GPU后，需要在CPU上加同步锁等待GPU运行完成(CPU给GPU发送一条指令，GPU就去算去了，等完了再返回给CPU，感觉就像是IO设备一样)

CUDA 为许多常用编程语言提供扩展，在此文以C++为例

.cu 是 CUDA 加速程序的文件扩展名

void CPUFunction()
{
  printf("This function is defined to run on the CPU.\n");
}

__global__ void GPUFunction()
{
  printf("This function is defined to run on the GPU.\n");
}

int main()
{
  CPUFunction();

  GPUFunction<<<1, 1>>>();
  cudaDeviceSynchronize();
}

• __global__ 关键字表明以下函数将在 GPU 上运行并可全局调用，而在此种情况下，则指由 CPU 或 GPU 调用。
• 通常，我们将在 CPU 上执行的代码称为主机代码，而将在 GPU 上运行的代码称为设备代码。
• 注意返回类型为 void。使用 __global__ 关键字定义的函数需要返回 void 类型。

 

GPUFunction<<<1, 1>>>();

• 通常，当调用要在 GPU 上运行的函数时，我们将此种函数称为已启动的核函数。
• 启动核函数时，我们必须提供执行配置，即在向核函数传递任何预期参数之前使用 <<< ... >>> 语法完成的配置。
• 在宏观层面，程序员可通过执行配置为核函数启动指定线程层次结构，从而定义线程组（称为线程块）的数量，以及要在每个线程块中执行的线程数量。稍后将在本实验深入探讨执行配置，但现在请注意正在使用包含 1 线程（第二个配置参数）的 1 线程块（第一个执行配置参数）启动核函数。

CUDA的函数在调用 的时候需要指定线程块数和线程数，这个地方我按照进程和线程理解的，前面一个参数定义进程数，后面一个定义每个进程有多少个线程，这就是CUDA并行原理

  cudaDeviceSynchronize();
为同步锁，CPU等待GPU运算完成

 

nvcc -arch=sm_70 -o out some-CUDA.cu -run

• nvcc 是使用 nvcc 编译器的命令行命令。

• 将 some-CUDA.cu 作为文件传递以进行编译。

• o 标志用于指定编译程序的输出文件。

• arch 标志表示该文件必须编译为哪个架构类型。本示例中，sm_70 将用于专门针对本实验运行的 Volta GPU 进行编译，但有意深究的用户可以参阅有关 arch 标志、虚拟架构特性 和 GPU特性 的文档。

• 为方便起见，提供 run 标志将执行已成功编译的二进制文件。

nvcc类似于gcc，一种编译器，需要指定架构，我估摸着GUDA针对不同架构做了不同的优化，-o输出文件

blockDim.x和threadDim.x为线程块数和每个线程块内线程数 ，可以类比进程与线程

一个经典的CUDA并行程序，每个线程只处理

threadIdx.x + blockIdx.x * blockDim.x及+stride的索引

Stride = gridDim.x * blockDim.x;

可以理解成输入有 N个数据，而CUDA有n个处理器，每个处理器仅处理第n,2n,...x*n(x*n<N)个数据

所以CUDA特别适合处理那些数据间没有相互依赖数据，比如图像，视频等等

CUDA的内存分配和释放使用

cudaMallocManaged，cudaFree，可以类比于malloc和free 

cudaGetLastError()，对于一些没有返回值的函数，当用户想知道该函数处理过程中是否出错就用该接口获取状态

cudaError_t err;
err = cudaMallocManaged(&a, N) 

常用的Err状态

 

最后看一个CUDA加法样例吧~

#include <stdio.h>
#include <assert.h>

inline cudaError_t checkCuda(cudaError_t result)
{
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
  return result;
}

void initWith(float num, float *a, int N)
{
  for(int i = 0; i < N; ++i)
  {
    a[i] = num;
  }
}

__global__
void addVectorsInto(float *result, float *a, float *b, int N)
{
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;

  for(int i = index; i < N; i += stride)
  {
    result[i] = a[i] + b[i];
  }
}

void checkElementsAre(float target, float *array, int N)
{
  for(int i = 0; i < N; i++)
  {
    if(array[i] != target)
    {
      printf("FAIL: array[%d] - %0.0f does not equal %0.0f\n", i, array[i], target);
      exit(1);
    }
  }
  printf("SUCCESS! All values added correctly.\n");
}

int main()
{
  const int N = 2<<20;
  size_t size = N * sizeof(float);

  float *a;
  float *b;
  float *c;

  checkCuda( cudaMallocManaged(&a, size) );
  checkCuda( cudaMallocManaged(&b, size) );
  checkCuda( cudaMallocManaged(&c, size) );

  initWith(3, a, N);
  initWith(4, b, N);
  initWith(0, c, N);

  size_t threadsPerBlock;
  size_t numberOfBlocks;

  threadsPerBlock = 256;
  numberOfBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;

  addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);

  checkCuda( cudaGetLastError() );
  checkCuda( cudaDeviceSynchronize() );

  checkElementsAre(7, c, N);

  checkCuda( cudaFree(a) );
  checkCuda( cudaFree(b) );
  checkCuda( cudaFree(c) );
}