CUDA笔记001：introduction

资源列表

https://zhuanlan.zhihu.com/p/346910129：整理好的资源列表

教程《CUDA C Programming Guide》(《CUDA C 编程指南》)：

• 官方英文：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#
• 中文翻译：https://github.com/HeKun-NVIDIA/CUDA-Programming-Guide-in-Chinese

本篇使用

正如同玄幻小说最初选择功法一样，一个合格的功法能够让你在前期少走很多弯路

此处暂时选择：https://zhuanlan.zhihu.com/p/34587739

代码可以自己独立写几个矩阵乘法加法或者参考sample，不在文档中留了

术语解读

异构计算（Heterogeneous Computing）：至少两种指令集架构（ISA）结合，如CPU+GPU，CPU+FPGA等。单纯一种CPU或GPU不算
https://en.wikipedia.org/wiki/Heterogeneous_computing

CUDA：(Compute Unified Device Architecture)
SM : Streaming Multiprocessor
SP : Streaming Processor
Warps： on a GPU are like threads on a CPU
Threads： on a GPU are like lanes on a SIMD CPU

kernel：在此处场景指的是Parallel GPU code

introduction

基本概念

此处，CPU + GPU的异构计算可以发挥功效：

• CPU：control flow，处理串行逻辑
• GPU：data flow，处理并行计算

CUDA所需要应对的场景，通常是数千个线程并行执行

为了能够实现更高强度的并行计算，GPU将更多的晶体管用于数据计算而不是数据缓存或控制流

编程实践

CUDA程序的源代码后缀为：.cu

核心概念：

• host：CPU及其内存
• device：指代GPU及其内存

典型的CUDA程序的执行流程如下：

1. 分配host内存，并进行数据初始化；
2. 分配device内存，并从host将数据拷贝到device上；
3. 调用CUDA的核函数在device上完成指定的运算；
4. 将device上的运算结果拷贝到host上；
5. 释放device和host上分配的内存。

简单来说，在CPU上进行复杂的控制，将计算任务抽象成单独模块交给GPU
由于是异构的，host代码和device代码被编译成不同架构的机器码，无法交换执行
本质上，可以把GPU看成CPU的一个加速设备。

Host 代码可以通过 API（如 CUDA、OpenCL）调用 Device 代码，实现数据传输和任务调度。

在CUDA编程中，主要的函数限定词有：

• __global__: 核函数，在device上执行，从host中调用，异步！！！！！
• __device__: 在device上执行，单仅可以从device中调用
• __host__（不写则默认）: 在host上执行，仅可以从host上调用

线程组织格式，grid和block都有对应的组织格式：

• 1D：向量计算
• 2D：图像处理，矩阵处理
• 3D：3D物理模拟，CT影像等

线程层级如下：

• grid：一个kernel只有一个grid
• block：一个block包括多个block
• thread：一个block包括多个thread

调用核函数时需要指定<<<grid, block>>>：

dim3 grid(8, 8, 4);   // 8×8×4 个 Block
dim3 block(16, 16, 4); // 每个 Block 内 16×16×4 线程

or

dim3 grid(8, 8);
dim3 block(16);

// 对于声明为dim3类型的，缺省位置为1，如grid的z轴

kernel_func<<<grid, block>>>(param);

由于SM的基本执行单元是包含32个线程的线程束，所以block大小一般要设置为32的倍数。

注意不同资源（memory，cache，controller）在不同计算层级的配置

注意，host和device两侧的内存管理是独立且隔离的，类似kernel中的userspace和kernel
需要使用cudaMemcpy来进行数据的移动，kind有：cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost，cudaMemcpyDeviceToDevice
cudaMemcpy()本身是同步的！，会阻塞CPU

但是，CUDA 6.0引入统一内存来共同管理host和device中的内存
cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flag=0);
但是，需要使用cudaDeviceSynchronize();手动同步

cuda async API：

• 流 (cudaStream_t)
• cudaMemcpyAsync()
  针对以上异步场景，需要考虑使用cudaDeviceSynchronize();在关键节点手动同步

由于host和device是异步的，因此等待GPU执行完成需要：

cudaDeviceReset(); // 释放 GPU 资源，重置 GPU 设备。
cudaDeviceSynchronize(); // 同步 CPU 和 GPU，等待所有 GPU 任务完成。

nvcc -arch=sm_75 -o test test.cu：使用nvcc编译（此处由于使用RTX2070，指定对应版本Compute Capability 7.5）

Q：不执行cudaDeviceSynchronize(); 会出现Segmentation fault (core dumped)，为什么？
A：CUDA kernel 是异步执行的，也就是说，kernel 启动后，CPU 代码不会等它执行完，而是直接往下执行。如果 CPU 在 GPU 计算完成之前访问 GPU 计算结果，就会导致未定义行为，可能会出现 Segmentation fault。