《CUDA C++ Programming Guide》阅读笔记（每页要点总结）

前言：零基础，之前跑过一点 python ，几乎没有搞过 C++ 。这次会利用 chatGPT 辅助阅读这本官方手册。学习 CUDA 编程。最终要达到啥效果呢，希望可以读懂 cudasift 的源码。

Chapter 1 使用GPU的好处

（3-4） cpu与gpu结构对比：gpu大部分晶体管都用于计算，而cpu会分心在缓存，控制等功能，以此来减少内存读取的时间花销。

Chapter 2 CUDA：通用的并行计算平台和编程模型

（5-6） cuda的应用：

• 库和中间件的：包括cudnn，tensorrt，cufft，cublas，curand，cusparse，cula之类的，还有一些没听过的
• 语言，c++，python之类的；
• 还有现实工业场景应用的。

Chapter 3 可伸缩的编程模型

（7） cuda的三个关键的概念

• 线程块层级
• 共享内存
• 障碍同步

问题不断拆分，在线程块中同步运行

（8） 一张线程块的结构图，Streaming Multiprocessors（SMs）提供了自动伸缩性，多核更快

（9） 整本书的结构

（11） 核函数基础概念，threadIdx，<<<线程块数量，每个线程块的线程数>>>

5.2 线程层级

（12-13） threadIdx是一个三维向量，线程块索引的计算，一个线程块因为内存的限制最多1024个线程，一个线程块可以是一维，二维，三维的（解答：https://www.zhihu.com/question/63319070）；在一个block内的线程可以通过共享内存和同步来统筹内存读取，通过在一个需要同步的点来调用__syncthreads()函数实现同步，这个函数就像一个障碍，所有block内的线程都要在这个障碍前停下来等待。

5.2.1 线程块簇

（14-15） 线程块簇

• 线程块簇的调度：线程块可以在流处理器（SMs）内调度，也可以GPU处理簇（GPC）内调度。
• 簇内线程块数：一个簇内最多就只能有8个线程块，可以通过cudaOccupancyMaxPotentialClusterSize这个api查一下最多支持多少个簇。
• 线程块簇的启动方式有2种：
  • 第一种：使用编译时核属性来修饰核函数， __cluster_dims__(X,Y,Z)
  • 第二种：cudaLaunchKernelEx

个人疑问：线程块簇有什么作用？

ChatGPT回答：在CUDA架构中，线程块（Thread Block）是一组并发执行的线程的集合。线程块通常是将任务分配给GPU的最小单位，且在一个多处理器上运行。然而，有时候一个多处理器可能无法容纳全部线程块，导致一些线程块无法立即执行。这时，Thread Block Clusters就发挥作用了。
Thread Block Clusters是多个线程块的组合，它们被调度在多个多处理器上执行。当一个多处理器上的线程块执行完成后，Thread Block Clusters可以通过调度器被分配到这个多处理器上继续执行，从而充分利用多处理器的计算资源。

也就是说，线程块集群的目的是提高并行性和资源利用率，通过在多个多处理器上同时执行线程块，提高CUDA程序的整体性能。

5.3 内存层级

（16） 线程与内存关系：

• 线程：每个线程都有私有的局部内存
• 线程块：每个线程块都一个块共享内存，对块内所有线程可见。
• 线程块簇：组成线程块簇的线程块，此时线程块内的共享内存变成了分布式的共享内存，对簇内的所有线程块可见。
• 额外的只读内存，内部的数据有更快的存取速度：
  • 常量存储（constant memory）：存储常量，使用__constant__关键字将数据声明为常量，
  • 纹理存储（texture memory）：存储图像或者二维数据，使用 texture 关键字声明为纹理

5.4 异构编程

（17） 异构编程的概念，gpu内存的allocation和deallocation，cpu与gpu之间数据传输。

统一内存（unified memory）：gpu和cpu公用一块内存地址空间，不需要分配拷贝之类的操作，host和device之间数据传输和访问更方便。

5.5 异步SIMT编程模型

（17） cuda编程模型提供异步编程模型来加速内存操作，包含：

• 线程之间的异步障碍
• 异步数据拷贝到GPU

5.5.1 异步操作

（17） 异步操作使用一个同步对象来同步操作的完成，怎么理解这句话呢，就是异步操作之后可能对执行别的操作，所以也需要在一个同步点停下来，确保所有的异步操作都完成了，才能进行下一步。

ChatGPT：同步对象可以是显式的同步原语，例如栅栏（barrier）或信号量（semaphore），也可以是由库提供的特定功能的对象，例如CUDA的cuda::barrier或cuda::pipeline。

线程范围（thread scope）：定义了同步对象可以同步异步操作的线程集合。

（18） 异构编程示意图

（19） 线程范围和定义

5.6 算力（compute capability）

（19） 也叫 SM版本。一般表示为 X.Y，X一样表示核心架构一样，Y只是表示新特性的更新。

6 编程接口

• C++拓展：定义自己的核函数，分配grid和block
• runtime库：
  • 在host运行的分配和解除分配device的内存的函数
  • 在host和device之间传输数据的函数
  • 管理多device的系统

runtime库是基于底层的

• C接口
• CUDA驱动接口：这个接口可以控制更为底层的概念，CUDA context（上下文，类似于cpu的进程），CUDA modules（模块，类似于cpu的动态加载库）

6.1 使用NVCC来进行编译

6.1.1 编译流程

离线编译
使用nvcc来编译，可以编译host和device的混合代码。但是两种代码在nvcc编译的时候会分开进行编译：

• 将device的代码编译成汇编形式的PTX或者二进制码