https://github.com/parallel101/course 前置条件 • 学过 C/C++ 语言编程。 • 理解 malloc/free 之类的概念。 • 熟悉 STL 中的容器、函数模板等。 • 英伟达 GTX900 及以上显卡。 • CUDA 11 及以上。 • CMake 3.18 及以上。 我负责监督你学习 第 0 章： Hello, world! CMake • 定义函数 kernel ，前面加上 __global__ 修 饰符，即可让他在 GPU 上执行。 • 不过调用 kernel 时，不能直接 kernel() ，而 是要用 kernel<<<1, 1>>>() 这样的三重尖括 号语法。为什么？这里面的两个 1 有什么用 ？稍后会说明。 • 运行以后，就会在 GPU 上执行 printf 了。 • 这里的 kernel 函数在 GPU GPU 上执行，称为核 函数，用 __global__ 修饰的就是核函数。 没有反应？同步一下！ • 然而如果直接编译运行刚刚那段代码，是不会打印出 Hello, world! 的。 • 这是因为 GPU 和 CPU 之间的通信，为了高效，是异 步的。也就是 CPU 调用 kernel<<<1, 1>>>() 后，并不 会立即在 GPU 上执行完毕，再返回。实际上只是把 kernel

，后半段主要介绍并行编程与优化。 1.课程安排与开发环境搭建： cmake 与 git 入门 2.现代 C++ 入门：常用 STL 容器， RAII 内存管理 3.现代 C++ 进阶：模板元编程与函数式编程 4.编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 通用寄存器： 32 位时代 • 32 位 x86 架构中的通用寄存器有： • eax, ecx, edx, ebx, esi, edi, esp, ebp • 其中 esp 是堆栈指针寄存器，和函数的调用与返回相关。 • 其中 eax 是用于保存返回值的寄存器。 通用寄存器： 64 位时代 • 64 位 x86 架构中的通用寄存器有： • rax, rcx, rdx, rbx, rsi 结论：如果发现编译器放弃了自动优化，可以 用 constexpr 函数迫使编译器进行常量折叠！ 不过， constexpr 函数中无法使用非 constexpr 的容器： vector, map, set, string 等 …… constexpr ：强迫编译器在编译期求值（续） 发现：会让编译变得很慢，因为这 50000 次迭代是在编译期进行的。 第 2 章：内联 调用外部函数： call 指令 @PLT 是

INode 一切节点的公共基类。 多态的经典案例 • IObject 具有一个 eatFood 纯虚函数，而 CatObject 和 DogObject 继承自 IObject ，他 们实现了 eatFood 这个虚函数，实现了多态。 • 注意这里解构函数（ ~IObject ）也需要是虚函数 ，否则以 IObject * 存储的指针在 delete 时只 会释放 IObject 里的成员，而不会释放 m_catFood 。所以 这里的解构函数也是多态的，他根据类型的不同 调用不同派生类的解构函数。 多态用于设计模式之“模板模式” • 这样之后如果有一个任务是要基于 eatFood 做文章，比如要重复 eatFood 两遍。 • 就可以封装到一个函数 eatTwice 里，这个函数只需接受他们共同的基类 IObject 作为参数，然后调 用 eatFood 这个虚函数来做事（而不是直接操作具体的猫和狗本身）。 dont-repeat-yourself ）， 也让函数的作者不必去关注点从猫和狗的其他具体细节，只需把握住他们统一具有的“吃”这个接口。 小知识： shared_ptr 如何深拷贝？ 浅拷贝： 深拷贝： 思考：能不能把拷贝构造函数也作为虚函数？ • 现在我们的需求有变，不是去对同一个对象调用两次 eatTwice ，而是先把对象复制一份 拷贝，然后对对象本身和他的拷贝都调用一次 eatFood 虚函数。 • 代码如下

数据。如果没有，则从内存中读取，并存储到缓存中； 如果有，则直接使用缓存中的数据。 • 这样一来，访问的数据量比较小时，就可以自动预先加 载到这个更高效的缓存里，然后再开始做运算，从而避 免从外部内存读写的超高延迟。 缓存的分级结构 查看高速缓存大小： lscpu • 可以看到我们 x86 电脑的缓存结构分为三级。 • 一级缓存分为数据缓存和指令缓存，其中数据缓存有 32 KB ， 数组就完全不会被读取，不会占用内 存带宽，从而带宽利用率是 100% ，因此比 AOS 快了 2 倍。 AOSOA ：两者得兼 • 还有一种办法就是让 MyClass 内部是 SOA ，而外部仍是一个 vector 的 AOS—— 这种内存布局称为 AOSOA 。 • 缺点是必须保证数量是 1024 的整数倍， 而且因为要两次指标索引，随机访问比较 烦。 字节都被读取到缓存了，但实际只用到了其中 4 字节，之 后又没用到剩下的 60 字节，导致浪费了 94% 的带宽。 • 虽说连续、顺序访问是最理想的，然而在使用哈希表等数据结构中，不 可避免的会通过哈希函数得到随机的地址来访问，且 Value 类型可能小 于 64 字节，浪费部分带宽。怎么办？ 解决：按 64 字节分块地随机访问 • 解决方法就是，把数据按 64 字节大小分块。随 机访问时，只随机块的位置，而块的内部仍然按

要访问一个多维数组，必须先创建一个表面对象 （ cudaSurfaceObject_t ）。 • 考虑到多维数组始终是需要通过表面对象来访问的，这 里我们让表面对象继承自多维数组。 • 在核函数中可以用 surf3Dread 和 surf3Dwrite 来读写 表面对象中的元素， x,y,z 参数指定要访问元素的坐标 ，要注意 x 必须乘以 sizeof( 元素类型 ) ，否则出错。 作为图形学专业硬件的能力。 CUDA 纹理对象：封装 • 表面对象访问数组是可读可写的。纹理对象也可以访问 数组，不过是只读的。好处是他可以通过浮点坐标来访 问，且提供了线性滤波的能力。 • 在核函数中可以通过 tex3D 来读取纹理中的值。 • 之所以纹理是因为 GPU 一开始是渲染图形的专用硬件 ，会用到一些贴图等，这就是二维的纹理。 • 当输入的浮点坐标不是整数时，由 GPU 硬件提供双线 和 resample_kernel 。 • 首先通过 advect_kernel 算出对流后要采样的位置，写入到 loc 。然后再对 clr 和 vel 分别从 loc 算出的位置重 新采样。核函数的 gridDim 通过上整除技巧保证每个元素都能访问到， blockDim 为 8x8x8=512 。 • 如果在 resample_kernel 需要读取 clr ，然后再写入 clr ，并且读写是不同的坐标位置。

.h> • 项目名 :: 函数名 (); 完整案例请看源码仓库： https://github.com/parallel101/course/tree/master/16/00 推荐的目录组织方式 • 头文件（项目名 /include/ 项目名 / 模块名 .h ）中写： • #pragma once • namespace 项目名 { • void 函数名 (); • } • 实现文件（项目名 实现文件（项目名 /src/ 模块名 .cpp ）中写： • #include < 项目名 / 模块名 .h> • namespace 项目名 { • void 函数名 () { 函数实现 } • } 完整案例请看源码仓库： https://github.com/parallel101/course/tree/master/16/00 推荐的目录组织方式 完整案例请看源码仓库： https://github 如果是一个类，则文件名应和类名相同，方便查找 （ Animal.cpp ）。 • 头文件中包含函数和类的声明，源文件则包含他们的实现。 七、只有头文件，没有源文件的情况 • 有时我们会直接把实现直接写在头文件里，这时可以没有与之对 应的源文件，只有一个头文件。 • 注意：在头文件里直接实现函数时，要加 static 或 inline 关键 字。 八、每新增一个功能模块，需要创建两个文件

OFF 防止不小心用了 GCC 才有的 特性。 • 此外，最好是在 project 指令前设置 CMAKE_CXX_STANDARD 这一系列变量，这样 CMake 可以在 project 函数里对编译器进行一些检测，看看他能不能支持 C++17 的特 性。 https://crascit.com/2015/03/28/enabling-cxx11-in-cmake/ 常见误区：小彭老师，我手动添加 "Visual Studio 2019", etc.) 一个标准的 CMakeLists.txt 模板 第 3 章：链接库文件 main.cpp 调用 mylib.cpp 里的 say_hello 函数 改进： mylib 作为一个静态库 改进： mylib 作为一个动态库 改进： mylib 作为一个对象库 https://www.scivision.dev/cmake-object-libraries/ C++ 特性等比较耗时，检测完会把结果存储到缓存中，这样第二遍运行 cmake -B build 时就可以直接用缓存的值，就不需要再检测一遍了。 如何清除缓存？删 build 大法了解一下 然而有时候外部的情况有所更新，这时候 CMake 里缓存的却是旧的值，会导致一系列问题。 这时我们需要清除缓存，最简单的办法就是删除 build 文件夹，然后重新运行 cmake -B build 。缓存是很多

实际上会被编译器翻译成他对应的 ASCII 码： 97 。写 ‘ a’ 和写 (char)97 是完全一样的，方便阅读的语法糖而已。 “char 即整数”思想应用举例 “char 即整数”思想应用举例 C 语言帮手函数 帮手函数大全 • isupper(c) 判断是否为大写字母（‘ A’ <= c && c <= ‘Z’ ）。 • islower(c) 判断是否为小写字母（‘ a’ <= c && c <= ‘z’ <= 31 或 c == 127 ）。 • toupper(c) 把小写字母转换为大写字母，如果不是则原封不动返回。 • tolower(c) 把大写字母转换为小写字母，如果不是则原封不动返回。 帮手函数大全 http://c.biancheng.net/ref/ctype_h/ 关于 char 类型的一个冷知识 • C 语言其实只规定了 unsigned char 是无符号 8 位整数， signed 。 • hello 每个字符都连续地排列在这个数组中，那么末尾的 0 是怎么回事？原来 C 语言的字符串因为只保留数组的 首地址指针（指向第一个字符的指针），在以 char * 类型 传递给其他函数时，其数组的长度无法知晓。为了确切知 道数组在什么地方结束，规定用 ASCII 码中的“空字符”也 就是 0 来表示数组的结尾。这样只需要一个首地址指针就 能表示一个动态长度的数组，高，实在是高。

，后半段主要介绍并行编程与优化。 1.课程安排与开发环境搭建： cmake 与 git 入门 2.现代 C++ 入门：常用 STL 容器， RAII 内存管理 3.现代 C++ 进阶：模板元编程与函数式编程 4.编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 for (auto _: bm) • 里面即可。他会自动决定要重复多少次， 保证结果是准确的，同时不浪费太多时间 。 运行结果 刚才的 BENCHMARK_MAIN 自动生成了一个 main 函数 ，从而生成一个可执行文件供你运行。运行后会得到测试 的结果打印在终端上。 命令行参数 他还接受一些命令行参数来控制测试的输出格式为 csv 等等，你可以调用 --help 查看更多用法。 CMake 用了工作窃取法来分配任务： 当一个线程 t1 做完自己队列里全部的工 作时，会从另一个工作中线程 t2 的队列 里取出任务，以免 t1 闲置浪费时间。 • 因此内部 for 循环有可能“窃取”到另一个 外部 for 循环的任务，从而导致 mutex 被重复上锁。 解决 1 ：用标准库的递归锁 std::recursive_mutex 解决 2 ：创建另一个任务域，这样不同域之间就不会窃取工作

协议标准化了本机代码调用 WebAssembly 和 WebAssembly 调用本机代码的通信 (messaging) 和 错误处理 (error handling) 。 什么是 waPC? WASM 外部函数 接口 FFI Input type Return type Platform 底层 I32 I32 Unknown waPC &[u8] Result,error