CUDA C++ 基础介绍

1. 异构编程

CPU+高速总线（PCIE/NVLink）+GPU（用于数学/科学计算，有自己的操作系统）

三步走：

1. 从 GPU 内存往 CPU 内存拷贝输入数据
2. 加载 GPU 代码并执行，在片上缓存数据来提升性能
3. 从 GPU 内存往 CPU 内存拷贝结果

并行计算：以 vector add 为例，计算 output 中第一个元素的过程和计算第二、三个元素的过程独立

2. GPU kernels: device code

函数定义

__global__ void mykernel(void) { }

• CUDA C++关键字/装饰器 global 表示这个函数运行在 GPU 上（给编译器的提示），通过 host code 或其它 device code 来调用
• NVCC：编译器。用于把源码划分为 host 和 device 部分
  • device 函数由 NVIDIA 编译器来执行
  • host 函数由 gcc，cl.exe 等标准 host 编译器执行

函数调用

mykernel<<<1, 1>>>();

• 三层尖括号表示调用 device code（也叫做 kernel launch），括号内的参数是 CUDA 核执行配置
• 第一个参数是 warp 数，第二个参数是每个 warp 包含的线程数

内存管理

• host 和 device 侧的内存是完全分隔的
  • device 侧的指针指向 GPU 内存，通常传递给 device code，不在 host 侧解引用
  • host 侧的指针指向 CPU 内存，通常不传给 device code，不在 device 侧解引用
  • 特例：pinned pointers，ATS，managed memory
• 用于 device 侧内存管理的 API（用法和 C 语言 API 类似）
  • cudaMalloc()
  • cudaFree()
  • cudaMemcpy()
• 这些 API 使用指针来分配、释放或者复制内存。指针只是一个数，它没有 metadata，不是正式的 C++ 对象。

3. 设备侧向量加法

__global__ void add(int *a, int *b, int *c){ c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x]; }

• grid 是线程和 block 的组合，是一种分层描述。grid - block（warp/warps) - threads
• blockIdx 是一个结构体/内置变量，有三个元素：.x，.y，.z。每个元素的索引从 0 开始，到 N-1 结束，其中 N 是 kernel launch 时传入的值。

#define N 512
{
     *a, *b, *c; 
     *d_a, *d_b, *d_c; 
     size = N * (); 
    (( **)&d_a, size);
    (( **)&d_b, size);
    (( **)&d_c, size);
    
    a = ( *)(size);
    (a, N);
    b = ( *)(size);
    (b, N);
    c = ( *)(size);
    
    (d_a, a, size, cudaMemcpyHostToDevice);
    
    (d_b, b, size, cudaMemcpyHostToDevice);
    
    add<<<N, >>>(d_a, d_b, d_c);
    
    (c, d_c, size, cudaMemcpyDeviceToHost);
    
    (a);
    (b);
    (c);
    (d_a);
    (d_b);
    (d_c);
     ;
}