表示回车（‘ \r’ ） • 27 表示 ESC 键（‘ \x1b’ ） • 127 表示 DEL 键（‘ \x7f’ ）等 • 0~31 和 127 这些整数，就构成了 ASCII 码中控制字符的部分。 关于控制字符的一个冷知识 • 在 Linux 命令行中启动 cat 。 • 试试按 Ctrl+R ， Ctrl+E ， Ctrl+C 等一系列 组合键，看到出现了什么？ • 可以看到显示的字符变成了 语言中规定字符类型为 char 类型，是个 8 位整数。 • 这是因为 ASCII 码只有 0~127 这些整数，而 8 位整数的表示范围是 2^8 也就是 0~255 ，足以表示所有 ASCII 字符了（多余的部分实际上被用于表示 中文）。 • char 和整数无异，例如 ‘ a’ 实际上会被编译器翻译成他对应的 ASCII 码： 97 。写 ‘ a’ 和写 (char)97 是完全一样的，方便阅读的语法糖而已。 的字符处写入 0 ，来提前结束字符串。例如在第 n 个字符写入 0 ，就会只保留前 n 个字符作为一个子字 符串，删除后半部分。 “0 结尾字符串”知识点应用举例 • C 语言所谓的字符串类型 char * 实际上就是个首地址指 针，如果让首地址指针向前移动 n 位，那就实现删除前 n 个字符的效果，而不用实际修改数组本身（更高效）。 C 语言转义符 • 常见的转义符： • ‘\n’ 换行符：另起一行（光标移到下一行行首）

移动的，我们仍可以通过移 动其指针的方式来实现双缓 冲（ std::swap ）。 对流部分 对流部分：计算对流后位置（ RK3 ） • 这里我参考了 Taichi 官方案例中的 stable_fluid.py 代码（二维定常流仿真），主要由 k-ye 编写 ，我学习 GAMES201 后贡献了支持 RK2 和 RK3 的版本。这里我们用高效的 CUDA 纹理对象 在 C++ 中重新实现了一遍，利用了硬件的三线性插值实现半拉格朗日（ 中重新实现了一遍，利用了硬件的三线性插值实现半拉格朗日（ semi-lagrangian ）对流。 对流部分：根据对流后位置重新采样 • 和 k-ye 思路不同的是我先在刚刚的 advect_kernel 算出对流后要采样的位置（ loc ），然 后再对 vel 和 clr 根据刚刚算得的 loc 移动位置。这样 RK3 的对流只需要算一遍，避免 重复对每个场都做一次对流的开销。 对流部分：最终实现 • 然后，在 SmokeSim::advection clr ，并且读写是不同的坐标位置。 • 因此对 clr 和 vel 使用了双缓冲，写入 clrNext 的同时读取 clr 没有冲突，写入完毕后对调 clrNext 和 clr 。 投影部分 投影部分 • 我们要模拟的流体是不可压缩的，因此有着无散度的特点： div v = 0 • 上式对时间求导，即 d(div v)/dt = div dv/dt = 0 ；带入 dv/dt = -p

unordered_map 中存储的表项数量，从而减轻哈 希的压力。但意味着键值在空间上需要具有一定的局域性，否 则 会浪费分块中一 部分空间。 然而我们这里是 要用他记录粒子 经过的点，因此 具有一定空间局 域性，能够被分 块优化。 实际上空间局域 性正是稀疏网格 能够实现的一大 前提，稍后详细 讨论。 在 16x16 分块的基础上，只用一个 bit 存储 图片解释稀疏的好处 传统稠密二维数组 unordered_map 就是充当根节点 (root node) 。 图片解释稀疏的好处 传统稠密二维数组 无边界稀疏分块哈希表 此外，还是按需分配内存，即使被写入的部分奇形怪状也不会浪费内存。 这些被写入的部分被称为激活元素 (active element) ，反之则是未激活 (inactive) 。 这就是稀疏的好处，按需分配，自动扩容。 分块则是利用了我们存储的数据常常有着空间局域性的特点，减轻哈希表的压 可以改成 a >> 3 。 >> 2 = 位运算 >> 对负数的处理 signed 类型的 >> n 会把最高位复制 n 次。 因为补码的特性，这导致负数 >> 的结果仍是负 数。 这样就实现了和 Python 一样的始终向下取整除 法。 >> 2 = unsigned 类型的位运算 >> 不一样 而 unsigned 类型的 >> n 会不会复制最高位， 只是单纯的位移，这会导致负数的符号位单独被位

com/parallel101/course 为什么往 int 数组里赋值 1 比赋值 0 慢一倍？ 第 1 章：内存带宽 cpu-bound 与 memory-bound • 通常来说，并行只能加速计算的部分，不能加速内存读写的部分 。 • 因此，对 fill 这种没有任何计算量，纯粹只有访存的循环体，并 行没有加速效果。称为内存瓶颈（ memory-bound ）。 • 而 sine 这种内部需要泰勒展开来计算，每次迭代计算量很大的 • 那么，频率相同的情况下，可以考虑插两块 8GB 的内存， 比插一块 16GB 的内存更快，不过价格可能还是翻倍的。 • 系统会自动在两者之间均匀分配内存，保证读写均匀分配 到两个内存上，实现内存的并行读写，这和磁盘 RAID 有 一定相似之处。 验证一下刚刚的 parallel_add 是不是用足了全部带宽 • 刚刚 a 数组的大小是 1024 MB 。 • 因为不光读取了 a 内存的空间局域性有关。 缓存行决定数据的粒度（续） • 所以我们设计数据结构时，应该把数据存 储的尽可能紧凑，不要松散排列。最好每 个缓存行里要么有数据，要么没数据，避 免读取缓存行时浪费一部分空间没用。 重新认识结构体 重新认识 AOS * * * * 重新认识 AOS *

新特性，都可以用。甚至可以把任何一个 C++ 项目的文件后缀名全部改成 .cu ，都能编 译出来。 • 这是 CUDA 的一大好处， CUDA 和 C++ 的关 系就像 C++ 和 C 的关系一样，大部分都兼容 ，因此能很方便地重用 C++ 现有的任何代码库 ，引用 C++ 头文件等。 • host 代码和 device 代码写在同一个文件内，这 是 OpenCL 做不到的。 编写一段在 GPU gcc 和 msvc ）生成 CPU 部分的指令码。然后送到真 正的 GPU 编译器生成 GPU 指令码。最后再链接成同一个文件 ，看起来好像只编译了一次一样，实际上你的代码会被预处理很 多次。 • 他在 GPU 编译模式下会定义 __CUDA_ARCH__ 这个宏，利用 #ifdef 判断该宏是否定义，就可以判断当前是否处于 GPU 模式 ，从而实现一个函数针对 GPU 和 CPU 生成两份源码级不同的 所以同步以后 CPU 也可以直接读取。 多个线程，并行地给数组赋值 • 刚刚的 for 循环是串行的，我们可以把线 程数量调为 n ，然后用 threadIdx.x 作为 i 索引。这样就实现了，每个线程负责给数 组中一个元素的赋值。 小技巧：网格跨步循环（ grid-stride loop ） • 无论调用者指定了多少个线程 （ blockDim ），都能自动根据给定的 n

来把数组设为空。 • void clear() noexcept; vector 容器： clear 配合 resize • resize 会保留原数组的前面部分不变，只 在后面填充上 0 。 • 如果需要把原数组前面的部分也填充上 0 ，可以先 clear 再 resize ，这是一个 常见的组合。 • void clear() noexcept; vector 容器： push_back back • 要注意的是 pop_back 函数的返回类型是 void ，也就是没有返回值，如果需要获取删除的 值，可以在 pop_back() 之前先通过 back() 获 取末尾元素的值，实现 pop 效果。 • a.back(); • 等价于： • a[a.size() - 1] • int &back() noexcept; • int const &back() const 函数返回之后解构） vector 容器： resize 到更大尺寸会导致 data 失效 • 当 resize 的目标长度大于原有的容量时， 就需要重新分配一段更大的连续内存，并 把原数组长度的部分移动过去，多出来的 部分则用 0 来填充。这就导致元素的地址 会有所改变，从而过去 data 返回的指针 以及所有的迭代器对象，都会失效。 vector 容器： resize 到更小尺寸不会导致 data

与 Python 和 C++ 不同， Java 放弃了花里胡哨的运算符重载，索性都采用成员函数 get put 来表示，非常明确。主要是为了把 get 和 put 作为接口函数，可以对应多个具体 实现。 错误示范 • 小彭老师说过，读取必须用 at 。 • 而这位同学却用了 [] 来读取 items 里的值。 • 乍看之下好像没错，运行结果也是正确的，但 这只是碰巧你的 items 里存在 structural-binding 很智能，刚刚说了 map 的类型是 pair 。 • 这个前者有一个 const 修饰，他就知道了，即使你这里是 auto & ，他对 K 部分也是会 变成 const K & 的。 • auto &[k, v] = tmp; • 等价于： • const K &k = tmp.first; • V &v = tmp.second; begin(); it != map.end(); ++it) 就一目了然了。 • for 里面第一部分，也就是初始化语句： it = map.begin() 代表从最左节点开始出发。 • 第二部分，也就是判断是否退出的条件： it != map.end() 判断是否抵达最右节点的下一个 。 • 第三部分，也就是每次循环后执行的更新语句： ++it 会让迭代器往下一个节点移动。 • 所以人话就是

GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 pbf 流体求解 12.C++ 在 ZENO 中的工程实践：从 primitive 说起 13.结业典礼：总结所学知识与优秀作业点评 I 硬件要求： 64 位（ 32 位时代过去了） 至少 2 为什么需要模板函数（ template ） • 避免重复写代码。 • 比如，利用重载实现“将一个数乘以 2” 这个 功能，需要： 为什么面向对象在 HPC 不如函数式和元编程香了？ 这个例子要是按传统的面向对象思想，可能是这样： 令 Int, Float, Double 继承 Numeric 接口类并实现 ，其中 multiply(int) 作为虚函数。然后定义： Numeric *twice(Numeric 来还不如重载方便了？ • 别担心， C++ 规定： • 当模板类型参数 T 作为函数参数时，则可 以省略该模板参数。自动根据调用者的参 数判断。 模板函数：特化的重载 • 有时候，一个统一的实现（比如 t * 2 ）满 足不了某些特殊情况。比如 std::string 就 不能用乘法来重复，这时候我们需要用 t + t 来替代，怎么办呢？ • 没关系，只需添加一个 twice(std::string)

GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 pbf 流体求解 12.C++ 在 ZENO 中的工程实践：从 primitive 说起 13.结业典礼：总结所学知识与优秀作业点评 I 硬件要求： 64 位（ 32 位时代过去了） 至少 2 可以容纳 4 个 float ，或 2 个 double 。 • 刚才的案例中只用到了 xmm 的低 32 位 用于存储 1 个 float 。 addss 是什么意思？ • 可以拆分成三个部分： add ， s ， s 1. add 表示执行加法操作。 2. 第一个 s 表示标量 (scalar) ，只对 xmm 的最低位进行运算；也可以是 p 表示矢量 (packed) ，一次对 xmm 为了效率我们可以尽量把常用函数定义在 头文件里，然后声明为 static 。这样调用 他们的时候编译器看得到他们的函数体， 从而有机会内联。 内联：当编译器看得到被调用函数（ other ）实现的时候 ，会直接把函数实现贴到调用他的函数（ func ）里。 局部可见函数： static 因为 static 声明表示不会暴露 other 给其 他文件，而且 func 也已经内联了 other ， 所以编译器干脆不定义