CUDA【2】Add

• 根据书中的第二个例子，我们需要实现两个一维数组的相加，为做对比，会先实现CPU版本，然后对应实现一边CUDA版本

Example2: add

• 这个是CPU版本，大体上可以用于了解基本功能，即数组的对应元素的相加然后用check求是否有误差

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

const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;

/*
    同样长度的一维数组之和
*/
void add(const double *x, const double *y, double *z, const int N);
void check(const double *z, const int N);

void add(const double *x, const double *y, double *z, const int N) {
  for (int n = 0; n < N; ++n) {
    z[n] = x[n] + y[n];
  }
}

void check(const double *z, const int N) {
  bool has_error = false;
  for (int n = 0; n < N; ++n) {
    if (fabs(z[n] - c) > EPSILON)
      has_error = true;
  }
  printf("%s\n", has_error ? "Has error!" : "No error!");
}

int main() {
  const int N = 10000;
  const int M = sizeof(double) * N;
  double *x = (double *)malloc(M);
  double *y = (double *)malloc(M);
  double *z = (double *)malloc(M);

  for (int n = 0; n < N; ++n) {
    x[n] = a;
    y[n] = b;
  }

  add(x, y, z, N);
  check(z, N);

  free(x);
  free(y);
  free(z);

  return 0;
}

• 这是CUDA版

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

const double EPSILON = 1.0e-3;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;

/*
    同样长度的一维数组之和
*/
void __global__ add(const double *x, const double *y, double *z);
void check(const double *z, const int N);

void __global__ add(const double *x, const double *y, double *z)
{
  const int n = blockDim.x * blockIdx.x + threadIdx.x;
  z[n] = x[n] + y[n];
}

void check(const double *z, const int N)
{
  bool has_error = false;
  for (int n = 0; n < N; ++n)
  {
    if (fabs(z[n] - c) > EPSILON)
      // printf("%f\n", z[n]);
      has_error = true;
  }
  printf("%s\n", has_error ? "Has error!" : "No error!");
}

int main()
{
  const int N = 100000;
  const int M = sizeof(double) * N;
  double *h_x = (double *)malloc(M);
  double *h_y = (double *)malloc(M);
  double *h_z = (double *)malloc(M);

  for (int n = 0; n < N; ++n)
  {
    h_x[n] = a;
    h_y[n] = b;
  }

  double *d_x, *d_y, *d_z;
  cudaMalloc((void **)&d_x, M);
  cudaMalloc((void **)&d_y, M);
  cudaMalloc((void **)&d_z, M);

  cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

  const int block_size = 160;
  const int grid_size = N / block_size;

  printf("%d\t%d\n", block_size, grid_size);

  // add<<<block_size, grid_size>>>(d_x, d_y, d_z);
  // add<<<block_size, grid_size>>>(d_x, d_y, d_z);
  
  // 这是为什么，当我将block_size，grid_size设置的比较大的时候就会有问题：P20
  add<<<block_size, grid_size>>>(d_x, d_y, d_z);
  // add<<<1, 1>>>(d_x, d_y, d_z);

  cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);

  check(h_z, N);

  free(h_x);
  free(h_y);
  free(h_z);

  cudaFree(d_x);
  cudaFree(d_y);
  cudaFree(d_z);
  return 0;
}

不同之处

• 相比于CPU版，我们需要在GPU中分配显存，函数原型是：
  • cudaError cudaMalloc(void **address, size_t size);
• 然后把内存中的数据拷贝到显存中去，即刚分配的显存的位置，函数原型是：
  • cudaError cudaMemcpy(void *dst, const void *src, size_t count, enum cudaMemcoyKind kind);
  • 其中枚举类别cudaMemcoyKind 包含：
    • cudaMemcpyHostToHost
    • cudaMemcpyHostToDevice
    • cudaMemcpyDeviceToHost
    • cudaMemcpyDeviceToDevice
    • cudaMemcpyDefault，根据指针所指的地址自动判断
• 释放资源使用cudaFree()

关于Kernel函数2

• 此前在HelloWorld中有对Kernel函数进行一些特点描述，此处会补充一些：
• 在Example2中我们通过块/线程号的计算来将数组相加的计算任务分配到CUDA的kernel中去，这种的任务分配个人感觉是比较核心的技能，后续可能得多关注这个。

核函数的要求

• 除了返回值为void，限定符为__global__之外，还要求：
  • 核函数还要求参数个数确定
  • 可以传递非指针变量
  • 且除非使用同一内存机制（如Jetson Nano）否则传入的指针必须指向Device内存
  • 核函数不能成为一个类的成员，通常做法是用包装函数调用核函数
  • 计算能力3.5之后引入动态并行，核函数内部可以调用其他核函数

核函数中if的引入

• 首先我们在分配块/线程的时候很有可能会出现需要的总线线程无法整除的情况，此时就很有必要引入if来判断一下是否到达临界点，然后及时退出

其他限定词

• 除了__global__之外（__global__表示核函数，一般主机调用，设备中执行）
• 还有__device__，表示设备函数，只能被其他核函数/设备函数调用，在设备中执行
• 以及__host__，就是普通的C++函数