r15b, r15w, r15d, r15 AT&T 汇编语言 GCC 编译器所生成的汇编语言就属于这种 返回值：通过 eax 传出 movl $42, %eax 相当于： eax = 42; 前 6 个参数：分别通过 edi ， esi ， edx ， ecx ， r8d ， r9d 传入 movl %edi, -4(%rsp) 相当于： *(rsp - 4) = edi; 开启优化： -O3 。 第 3 章：指针 编译器傻了吗？ 为什么编译器不优化掉 *c = *a ？ 指针别名现象（ pointer aliasing ） 考虑这种调用方式： b 和 c 指向同一个变量。 优化前相当于： b = a; b = b; 最后 b 变成了 a 。 如果优化了： b = b; 最后 b 没有改变。 导致优化后结果不一样，这就是 编译器放弃优化的原因。 告诉编译器别怕指针别名： __restrict 个 int 的加法，从而更加高 效但这样有个缺点，那就是数组的大小必须为 4 的整数倍 否则就会写入越界的地址！ 如果不是 4 的倍数？边界特判法 看不懂？很简单，假设 n = 1023 ： 先对前 1020 个元素用 SIMD 指令填入，每次处理 4 个 剩下 3 个元素用传统的标量方式填入，每次处理 1 个 思想：对边界特殊处理，而对大部分数据能够矢量化 编译器做优化时会自动处理边界特

的执行队列上，然后立即 返回，并不会等待执行完毕。 • 因此可以调用 cudaDeviceSynchronize() ，让 CPU 陷 入等待，等 GPU 完成队列的所有任务后再返回。从而 能够在 main 退出前等到 kernel 在 GPU 上执行完。 定义在 GPU 上的设备函数 • __global__ 用于定义核函数，他在 GPU 上执行，从 CPU 端通过三重尖括号语法调 用，可以有参数，不可以有返回值。 如需总的线程编号： blockDim * blockIdx + threadIdx 分离 __device__ 函数的声明和定义：出错 • 默认情况下 GPU 函数必须定义在同一个文件里。 如果你试图分离声明和定义，调用另一个文件里 的 __device__ 或 __global__ 函数，就会出错 。 分离 __device__ 函数的声明和定义：解决 • 开启 CMAKE_CUDA_ CMAKE_CUDA_SEPARABLE_COMPILATION 选 项（设为 ON ），即可启用分离声明和定义的支持。 • 不过我还是建议把要相互调用的 __device__ 函数放在 同一个文件，这样方便编译器自动内联优化（第四课讲 过）。 两种开启方式：全局有效 or 仅针对单个程序 只对 main 这个程序启用： 对下方所有的程序启用（推荐）： 顺便一提， CXX_STANDARD 和 CUDA_ARCHITECTURES

里，然后在需要用到 hello() 这个声 明的地方，打上一个记号， #include “hello.h” 。然后用一个小程序，自动在编译前把引号 内的文件名 hello.h 的内容插入到记号所在的位置，这样不就只用编辑 hello.h 一次了嘛 ~ • 后来，这个编译前替换的步骤逐渐变成编译器的了一部分，称为预处理阶段， #define 定 义的宏也是这个阶段处理的。 • 此外，在实现的文件 hello hello.cpp 被修改时，比如改成 hello(int) ，编译器能够发现 hello.h 声明的 hello() 和定 义的 hello(int) 不一样，避免“沉默的错误”。 编译前替换 编译前替换 头文件 - 批量插入几行代码的硬核方式 • 实际上 cstdio 也无非是提供了 printf 等一系列函数声明的头文件而已，实际的实现是在 libc.so 这个动态库里。其中 hello() 和定义的 hello(int) 不一样，避免“沉默的错误”（虽然对支持重载的 C++ 不奏 效） 2. 可以让 hello.cpp 中的函数需要相互引用时，不需要关心定义的顺序。 编译前替换 编译前替换 头文件进阶 - 递归地使用头文件 • 在 C++ 中常常用到很多的类，和函数一样，类的声明也会被放到头文件中。 • 有时候我们的函数声明需要使用到某些类，就需要用到声明了该类的头文件，像这样递归

non-constexpr 函数。而且 constexpr 函数必须是内联 （ inline ）的，不能分离声明和定义在另一个文件里。标准库的很多函数如 std::min 也是 constexpr 函数，可以放心大胆在模板尖括号内使用。 模板的难题：移到另一个文件中定义 • 如果我们试着像传统函数那样分离模板函数的声明与实现： • 就会出现 undefined reference 错误： 模板的难题：移到另一个文件中定义（续） 里只看到 sumto<> 函数的两份声明，从而出错。 • 解决：在看得见 sumto<> 定义的 sumto.cpp 里，增加两个显式编译模板的声明： 一般来说，我会建议模板不要 分离声明和定义，直接写在头 文件里即可。如果分离还要罗 列出所有模板参数的排列组合 ，违背了开 - 闭原则。 模板的惰性：延迟编译 • 要证明模板的惰性，只需看这个例子： • 要是编译器哪怕细看了一眼：字符串怎么可能被写入呢？肯定是会出错的。 1. 当函数有多条 return 语句时，所有语句的返回类型必须一致，否则 auto 会报错。 2. 当函数没有 return 语句时， auto 会被推导为 void 。 3. 如果声明和实现分离了，则不能声明为 auto 。比如： auto func(); // 错误 C++ 特性：引用（ int & ） • 众所周知， C++ 中有一种特殊的类型，叫做引用。只需要在原类型后面加一个 &

const &val); vector 容器： resize • 调用 resize(n) 的时候，如果数组里面不足 n 个元素，假设是 m 个，则他只会用 0 填充新增 的 n - m 个元素，前 m 个元素会保持不变。 • vector a = {1, 2}; • a.resize(4); • 等价于： • vector a = {1, 2, 0, 0}; • void resize(size_t n); vector 容器： resize • 调用 resize(n) 的时候，如果数组已有超过 n 个元素，假设是 m 个，则他会删除多出来的 m - n 个元素，前 n 个元素会保持不变。 • vector a = {1, 2, 3, 4, 5, 6}; • a.resize(4); • 等价于： • vector a = {1, 2 容器： resize • 调用第二个重载 resize(n, val) 的时候，如果数组 里面不足 n 个元素，假设是 m 个，则他只会用第 二个参数 val 填充新增的 n - m 个元素，前 m 个 元素会保持不变。 • vector a = {1, 2}; • a.resize(4, 233); • 等价于： • vector a = {1, 2, 233

应用程序中很常见，比如浏览 器在后台下载文件的同时，用户仍然可以 用鼠标操作其 UI 界面。 没有多线程：程序未响应 • 没有多线程的话，就必须等文件下载完了 才能继续和用户交互。 • 下载完成前，整个界面都会处于“未响应”状 态，用户想做别的事情就做不了。 现代 C++ 中的多线程： std::thread • C++11 开始，为多线程提供了语言级别的 支持。他用 std::thread download 函数才会出师未捷身先死 ——还没开始执行他的线程就被销毁了。 解构函数不再销毁线程： t1.detach() • 解决方案：调用成员函数 detach() 分离该 线程——意味着线程的生命周期不再由当 前 std::thread 对象管理，而是在线程退 出以后自动销毁自己。 • 不过这样还是会在进程退出时候自动退出 。 解构函数不再销毁线程：移动到全局线程池 • 但是 在运行物理解算的时候，界面会卡住，算完一帧后窗口才能刷新一遍 ，导致解算过程中基本别想做事，这一定程度上归功于 opengl 原始的单线程设计。 • 正面教材： zeno 可以在解算过程中，随时拖动滑块看前几帧的结果，编辑场景图，修改 节点间的连接，为下一次解算做准备，同时当前已经启动的物理解算还能在后台继续正常 运行。虽然 zeno 也用了 opengl ，但他用多进程成功在 opengl 的百般拖后腿下实现了

getter/setter 封装。 • 各个成员之间相互正交，比如数学矢量类 Vec3 ，就没必要去搞封装，只会让程序员 变得痛苦，同时还有一定性能损失：特别 是如果 getter/setter 函数分离了声明和定 义，实现在另一个文件时！ C++ 思想： RAII （ Resource Acquisition Is Initialization ） 资源获取视为初始化，反之，资源释放视为销毁 制权后，仍希望访问到对象的需求。 • 如果还是用 p 去访问的话，因为被移动构 造函数转移了， p 已经变成空指针，从而 出错。 解决方案：提前获取原始指针 • 最简单的办法是，在移交控制权给 func 前，提前通过 p.get() 获取原始指针： 解决方案：提前获取原始指针（续） • 不过你得保证 raw_p 的存在时间不超过 p 的生命周期，否则会出现危险的空悬指 针。比如右边这样： 更智能的指针： 根据实际情况 （主要是看所有权），判断要用哪一种智能指针： 1. unique_ptr ：当该对象仅仅属于我时。比如：父窗口中指向子 窗口的指针。 2. 原始指针：当该对象不属于我，但他释放前我必然被释放时。 有一定风险。比如：子窗口中指向父窗口的指针。 3. shared_ptr ：当该对象由多个对象共享时，或虽然该对象仅仅 属于我，但有使用 weak_ptr 的需要。 4. weak_ptr

结构体的内存布局： AOS 与 SOA • AOS （ Array of Struct ）单个对象的属性紧挨着存 • xyzxyzxyzxyz • SOA （ Struct of Array ）属性分离存储在多个数组 • xxxxyyyyzzzz • AOS 必须对齐到 2 的幂才高效， SOA 就不需要。 • AOS 符合直觉，不一定要存储在数组这种线性结构， 而 SOA 可能无法保证多个数组大小一致。 解决方案就是，把分块的大小调的更大一些，比 如 4KB 那么大，即 64 个缓存行，而不是一个。 • 这样一次随机访问之后会伴随着 64 次顺序访问， 能被 CPU 检测到，从而启动缓存行预取，避免了 等待数据抵达前空转浪费时间。 页对齐的重要性 • 为什么要 4KB ？原来现在操作系统管理内存是用分页 （ page ），程序的内存是一页一页贴在地址空间中的， 有些地方可能不可访问，或者还没有分配，则把这个页设 以后只分配了 4KB 的内存。等到用户访问了 a[1024] ，也就是触及了下一个页面，他才 会继续分配一个 4KB 的页面，这时才 8KB 被实际分配。比如这里我们分配了 16GB 内 存，但是只访问了他的前 4KB ，这样只有一个页被分配，所以非常快。 • 其实操作系统惰性分配的特性，也是 SPGrid （ Sparsely-Paged-Grid ）得以实现的基础 ，他利用 mmap 分配比机器大得多的内存（比如

疏网格、位运算、浮点的二进制格式、内存带宽优 化 面向人群：图形学、 CFD 仿真、深度学习编程人 员 第 0 章：稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 来存储 分离 read/write/create 三种访问模式 foreach 直接给出当前坐标指向的值 改用 unordered_map 来存储 unordered_map 手动 read(i, j) 也一样速度 以后只分配了 4KB 的内存。等到用户访问了 a[1024] ，也就是触及了下一个页面，他才 会继续分配一个 4KB 的页面，这时才 8KB 被实际分配。比如这里我们分配了 16GB 内 存，但是只访问了他的前 4KB ，这样只有一个页被分配，所以非常快。 实验：那如果分配超过机器内存容量的空间会怎样 • 既然是操作系统的内存是惰性分配给用户程 序的，分块大小就是 4KB ，那么是不是可 以利用这一点实现稀疏？

注意不论是项目自己的头文件还是外部的系统的头文件，请全部统一采用 < 项目名 / 模块名 .h> 的格式。不要用 “模块名 .h” 这种相对路径的格式，避 免模块名和系统已有头文件名冲突。 十、依赖其他模块但不解引用，则可以只前向声明不导入头文件 • 如果模块 Carer 的头文件 Carer.h 虽然引用了其他模块中的 Animal 类，但 是他里面并没有解引用 Animal ，只有源文件 Carer.cpp 解引用了 位于 bin 目录，而不是 lib 目录，这是为什么呢？ • 因为 Windows 要求 exe 和 dll 位于同一目录，否则 exe 在运行时就会找不到 dll 。 • 为了符合 Linux 分离 bin 和 lib 的组织格式，又要伺候 Windows 的沙雕同目录规则，我们通 常把 dll 动态库文件视为“可执行文件”和 exe 一起放到 bin 目录，而静态库则没有运行时必须 同目录的限制，所以可以照常放到