CUDA 编程

CUDA 一种异构计算平台

CUDA 是 NVIDIA 推出的用于自家 GPU 的并行计算框架，也就是说 CUDA 只能在 NVIDIA 的 GPU 上运行，而且只有当要解决的计算问题是可以大量并行计算的时候才能发挥 CUDA 的作用。CUDA 的主要作用是连接 GPU 和应用程序，方便用户通过 CUDA 的 API 调度 GPU 进行计算。

一个CUDA应用通常可以分解为两部分，

CPU 主机端代码
GPU 设备端代码

CUDA nvcc 编译器会自动分离你代码里面的不同部分，host 代码用 cpp 写成，使用本地的 g++ 编译器编译，设备端代码，也就是核函数，用 CUDA C 编写，通过 nvcc 编译，链接阶段，在内核程序调用或者明显的 GPU 设备操作时，添加运行时库。

CUDA 编程

简单示例代码

hello world 例子：

/*
*hello_world.cu
*/
#include<stdio.h>
__global__ void hello_world(void) {
  printf("GPU: Hello world! \n");
}

int main(int argc, char **argv) {
  printf("CPU: Hello world!\n");
  hello_world<<<1,10>>>();
  cudaDeviceReset(); //if no this line, it can not output hello world from gpu
  return 0;
}

简单介绍其中几个关键字

1	__global__ // 是告诉编译器这个是个可以在设备上执行的核函数

1	hello_world<<<1, 10>>>(); // 其中变量的含义是<<<线程块的个数,每个线程块中线程的个数>>> 一个核函数被执行的次数就是两个参数的乘积

1	cudaDeviceReset();

这句话如果没有，则不能正常的运行，因为这句话包含了隐式同步，GPU 和 CPU 执行程序是异步的，核函数调用后成立刻会到主机线程继续，而不管 GPU 端核函数是否执行完毕，所以上面的程序就是GPU 刚开始执行，CPU已经退出程序了，所以我们要等 GPU 执行完了，再退出主机线程。

调用核函数

核函数就是在 CUDA 模型上诸多线程中运行的那段串行代码，这段代码在 GPU 上运行，用 NVCC 编译，产生的机器码是 GPU 的机器码，所以我们写 CUDA 程序就是写核函数，第一步我们要确保核函数能正确的运行产生正确的结果，第二优化 CUDA 程序的部分，无论是优化算法，还是调整内存结构，线程结构都是要调整核函数内的代码，来完成这些优化的。
我们一直把我们的 CPU 当做一个控制者，运行核函数，要从 CPU 发起。

1	kernel_name<<<grid, block, share_mem, stream>>>(argument list);

<<<grid, block, share_mem, stream>>> 是对 GPU 代码执行的线程结构的配置。我们通过 CUDA C 内置的数据类型 dim3 类型的变量来配置 grid 和 block。

grid: grid 中 block 的个数
block: 每个 block 中 thread 的布局
是一个可选参数，用于设置每个 block 除了静态分配的 shared memory 外，最多能动态分配的 shared memory 大小，单位为字节，默认为 0。
是一个可选参数，是 cudaStream_t 类型，初始值为 0，用于表示该核函数处于哪个流中。

例如：

1	kernel_name<<<4, 8>>>(argument list);

表现为

可以用 threadIdx.x 和 blockIdx.x （dim3 类型，可以为x, y, z）来组合获得对应的线程的唯一标识。当主机启动了核函数，控制权马上回到主机（不阻塞），而不是主机等待设备完成核函数的运行。想要主机等待设备端执行可以用下面这个指令：

1	cudaError_t cudaDeviceSynchronize(void);

当然，有些操作要阻塞，比如内存拷贝，因为要用到 host。

编写核函数

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
  int i = threadIdx.x;
  C[i] = A[i] + B[i];
}

限定符	执行	调用	备注
__global__	设备端执行	可以从主机调用也可以从计算能力3以上的设备调用	必须有一个void的返回类型
__device__	设备端执行	设备端调用
__host__	主机端执行	主机调用	可以省略

Kernel核函数编写有以下限制

只能访问设备内存
必须有 void 返回类型
不支持可变数量的参数
不支持静态变量
显示异步行为

完整示例代码

一般 CUDA 程序分成下面这些步骤：

分配 GPU 内存
拷贝内存到设备
调用 CUDA 内核函数来执行计算
把计算完成数据拷贝回主机端
内存释放

#include <cuda_runtime.h>
#include <stdio.h>
// 错误检验的宏
#define CHECK(call)\
{\
  const cudaError_t error=call;\
  if(error!=cudaSuccess)\
  {\
      printf("ERROR: %s:%d,",__FILE__,__LINE__);\
      printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
      exit(1);\
  }\
}

__global__ void sumArraysGPU(float* a, float* b, float* res) {
  int i = threadIdx.x;
  res[i] = a[i] + b[i];
}

void initialData(float* vec, int n) {
    for (int i = 0; i < n; i++) {
        vec[i] = (float)i;
    }
}

int main(int argc,char **argv) {
  int dev = 0;
  cudaSetDevice(dev);

  int nElem = 32;
  printf("Vector size:%d\n", nElem);
  int nByte = sizeof(float)*nElem;
  float *a_h = (float*)malloc(nByte); // 输入数据 a
  float *b_h = (float*)malloc(nByte); // 输入数据 b
  float *res_from_gpu_h = (float*)malloc(nByte); // 用于接受从 gpu 返回的结果
  memset(res_from_gpu_h, 0, nByte);

  float *a_d, *b_d, *res_d; // 核函数的输入核输出地址（在 gpu 上申请的内存）
  CHECK(cudaMalloc((float**)&a_d, nByte)); // 在 gpu 上申请内存
  CHECK(cudaMalloc((float**)&b_d, nByte));
  CHECK(cudaMalloc((float**)&res_d, nByte));

  initialData(a_h, nElem);
  initialData(b_h, nElem);

  CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice)); // 将输入数据拷贝到 gpu
  CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

  dim3 block(nElem);
  dim3 grid(nElem / block.x);
  sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d); // 调用核函数
  printf("Execution configuration<<<%d, %d>>>\n", block.x, grid.x);

  CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost)); // 将核函数结果从 gpu 返回给 主机

  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);

  free(a_h);
  free(b_h);
  free(res_from_gpu_h);

  return 0;
}

1	nvcc xxx.cu -o a.out

caijizhuo

cuda编程

CUDA 编程

CUDA 一种异构计算平台

CUDA 编程

简单示例代码

调用核函数

编写核函数

完整示例代码