视频链接：【NVIDIA】CUDA官方入门课 - 01

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
int main(void){
  int *a, *b, *c;  // host copies of a, b, c
  int *d_a, *d_b, *d_c;  // device copies of a, b, c
  int size = N * sizeof(int);
  // Alloc space for device dopies of a, b, c
  cudaMalloc((void **)&d_a, size);
  cudaMalloc((void **)&d_b, size);
  cudaMalloc((void **)&d_c, size);
  // Alloc space for host copies of a, b, c and setup input values
  a = (int *)malloc(size); random_ints(a, N);
  b = (int *)malloc(size); random_ints(b, N);
  c = (int *)malloc(size);
  
  // Copy inputs to device
  cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);  // 可以不加cudaMemcpyHostToDevice，但是加上相当于提供了一层校验
  cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
  // Launch add() kernel on GPU with N blocks (N copies of add())
  add<<<N,1>>>(d_a, d_b, d_c);
  
  // Copy result back to host
  cudaMemcpy(c, d_c, size cudaMemcpyDeviceToHost);
  
  // Cleanup
  free(a); free(b); free(c);
  cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
  return 0;
}

• cuda开头的API：cuda runtime API。如果执行错误，会返回错误码
• cuda符合C++ 2014标准，但是不支持C++标准库

4. CUDA线程

• 术语：一个block可以被划分为并行线程

• 改写add()来使用并行线程，而非并行blocks

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

  使用threadIdx.x代替blockIdx.x

  main函数调用该kernel时，也需要修改调用参数

  add<<<1, N>>>();
  

5. 线程块和线程

• 如果每个block有M个线程，则 int index = threadIdx.x + blockIdx.x * M;

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

#define N (2048*2048)
#define THREAD_PRE_BLOCK 512
...
add<<<N/THREAD_PRE_BLOCK, THREAD_PRE_BLOCK>>>(d_a, d_b, d_c);
...

6. 处理任意大小的向量

__global__ void add(int *a, int *b, int *c, int n){
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  if (index < n){
    c[index] = a[index] + b[index];
  }
}

add<<<(N + M - 1) / M, M>>>(d_a, d_b, d_c, N);

相同块内的线程可以互相沟通和同步。但是不同块中的线程没有这个能力。

Q&A

• cudaMemcpy通过高速总线实现数据传输。会有传输耗时，它的带宽比GPU内部带宽要小，会限制GPU的性能。因此尽量减少host和device侧的数据传输。
• host和device侧的size可以确保是一致的，不会有LP/LLP差异，如host侧的long是64bit，那么device侧的long也是64bit。但在host侧使用的编译器类型受cuda限制。
• grid的大小通常是根据问题规模来确定的
• threadIdx<1024，没有其他限制
• host和device侧的代码可以放在两个文件中，只要能用C/C++方式编译起来