深入浅出访存优化 by 彭于斌（ @archibate ） 往期录播： https://www.bilibili.com/video/BV1fa411r7zp 课程 PPT 和代码： https://github.com/parallel101/course 为什么往 int 数组里赋值 1 比赋值 0 慢一倍？ 第 1 章：内存带宽 cpu-bound 与 memory-bound • 通常来说，并行只能加速计算的部分，不能加速内存读写的部分 。 • 因此，对 fill 这种没有任何计算量，纯粹只有访存的循环体，并 行没有加速效果。称为内存瓶颈（ memory-bound ）。 • 而 sine 这种内部需要泰勒展开来计算，每次迭代计算量很大的 循环体，并行才有较好的加速效果。称为计算瓶颈（ cpu- bound ）。 • 并行能减轻计算瓶颈，但不减轻内存瓶颈，故后者是优化的重点 并行能减轻计算瓶颈，但不减轻内存瓶颈，故后者是优化的重点 。 浮点加法的计算量 • 冷知识：并行地给浮点数组每个元素做一次加法反而更慢。 • 因为一次浮点加法的计算量和访存的超高延迟相比实在太少了。 • 计算太简单，数据量又大，并行只带来了多线程调度的额外开销 。 • 小彭老师经验公式： 1 次浮点读写 ≈ 8 次浮点加法 • 如果矢量化成功（ SSE ）： 1 次浮点读写 ≈ 32 次浮点加法 • 如果 CPU 有

都可以调用 • 这样相当于把 constexpr 函数自动变成修饰 __host__ __device__ ，从而两边都可以调用。 • 因为 constexpr 通常都是一些可以内联的函数，数学计 算表达式之类的，一个个加上太累了，所以产生了这个 需求。 • 不过必须指定 --expt-relaxed-constexpr 这个选项才能 用这个特性，我们可以用 CMake 的生成器表达式来实 如需总的线程编号： blockDim * blockIdx + threadIdx 分离 __device__ 函数的声明和定义：出错 • 默认情况下 GPU 函数必须定义在同一个文件里。 如果你试图分离声明和定义，调用另一个文件里 的 __device__ 或 __global__ 函数，就会出错 。 分离 __device__ 函数的声明和定义：解决 • 开启 CMAKE_CUDA_ CMAKE_CUDA_SEPARABLE_COMPILATION 选 项（设为 ON ），即可启用分离声明和定义的支持。 • 不过我还是建议把要相互调用的 __device__ 函数放在 同一个文件，这样方便编译器自动内联优化（第四课讲 过）。 两种开启方式：全局有效 or 仅针对单个程序 只对 main 这个程序启用： 对下方所有的程序启用（推荐）： 顺便一提， CXX_STANDARD 和 CUDA_ARCHITECTURES

cudaMalloc3DArray 用于分配一个三维数组。 各维度上的大小通过 cudaExtent 指定，方 便起见我们的 C++ 封装类用了 uint3 表示 大小。 • GPU 的多维数组有特殊的数据排布来保障 访存的高效，和我们 CPU 那样简单地行主 序或列主序（如 a[x + nx * y] ）的多维数组 不一样。 • 随后可用 cudaMemcpy3D 在 GPU 的三 维数组和 CPU 的三维数组之间拷贝数据。 ， n+100 钳制到 n-1 。 • cudaBoundaryModeZero ：对于读来说越界会读取到 0 ；对于写来说越界会放弃写入，不修改数组中的任 何值。 • 表面对象保障了高效的访存，并且自动判断越界，体 现了 GPU 作为图形学专业硬件的能力。 CUDA 纹理对象：封装 • 表面对象访问数组是可读可写的。纹理对象也可以访问 数组，不过是只读的。好处是他可以通过浮点坐标来访 和 k-ye 思路不同的是我先在刚刚的 advect_kernel 算出对流后要采样的位置（ loc ），然 后再对 vel 和 clr 根据刚刚算得的 loc 移动位置。这样 RK3 的对流只需要算一遍，避免 重复对每个场都做一次对流的开销。 对流部分：最终实现 • 然后，在 SmokeSim::advection 中调用 advect_kernel 和 resample_kernel 。

进阶：模板元编程与函数式编程 4.编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 。 • 所以， download 函数才会出师未捷身先死 ——还没开始执行他的线程就被销毁了。 解构函数不再销毁线程： t1.detach() • 解决方案：调用成员函数 detach() 分离该 线程——意味着线程的生命周期不再由当 前 std::thread 对象管理，而是在线程退 出以后自动销毁自己。 • 不过这样还是会在进程退出时候自动退出 。 解构函数不再销毁线程：移动到全局线程池 。 • 反面教材： blender 在运行物理解算的时候，界面会卡住，算完一帧后窗口才能刷新一遍 ，导致解算过程中基本别想做事，这一定程度上归功于 opengl 原始的单线程设计。 • 正面教材： zeno 可以在解算过程中，随时拖动滑块看前几帧的结果，编辑场景图，修改 节点间的连接，为下一次解算做准备，同时当前已经启动的物理解算还能在后台继续正常 运行。虽然 zeno 也用了 opengl

进阶：模板元编程与函数式编程 4.编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 getter/setter 封装。 • 各个成员之间相互正交，比如数学矢量类 Vec3 ，就没必要去搞封装，只会让程序员 变得痛苦，同时还有一定性能损失：特别 是如果 getter/setter 函数分离了声明和定 义，实现在另一个文件时！ C++ 思想： RAII （ Resource Acquisition Is Initialization ） 资源获取视为初始化，反之，资源释放视为销毁 这个编译器自动生成的初始化列表构造函 数，除了可以指定全部成员来构造以外， 还可以指定部分的成员，剩余没指定的保 持默认。 • 不过你得保证那个没指定的有在类成员定 义里写明 {} 初始化，否则有可能会变成内 存里的随机值。 • 顺便一提， C++20 中还可以通过指定名称来跳顺序： 编译器默认生成的构造函数：初始化列表（妙用，解决函数多返回值） • 典型的例子包括，图形学某知名应用中， 可以简化函数具有多个返回值的处理。

疏网格、位运算、浮点的二进制格式、内存带宽优 化 面向人群：图形学、 CFD 仿真、深度学习编程人 员 第 0 章：稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 来存储 分离 read/write/create 三种访问模式 foreach 直接给出当前坐标指向的值 改用 unordered_map 来存储 unordered_map 手动 read(i, j) 也一样速度 i % 2 的计算时间，完全隐藏在内存 的超高延迟里了。 • 可见，当数据量足够大，计算量却不多时，读写 数据量的大小唯一决定着你的性能。 • 特别是并行以后，计算量可以被并行加速，而访 存却不行。 使用 int8_t ：每个占据 1 字节 • 因此我们可以把数据类型变小，这样所需的内存 量就变小，从而内存带宽也可以减小！ • 对于右边这种内存瓶颈的循环体，从 4 字节的 int 试图解决：用 uint8_t 表示，定点数系数调小到 2 • 注意到我们的值始终是正数，因此可以用无符号的 uint8_t （可以容纳 0 到 255 ），然后把刚刚的系数 100 改小到 2 ，成功算对结果了，代价是精度损失了 不少。 • 其实 GPU 存储贴图一般也是用的定点数 uint8_t （范围从 0 到 255 ），着色器在读取的时候才会把他 转换成 float （范围从 0.0

之间的区别）。在编写 C++ 时，也应该尽可能 的避免使用诸如 void* 之类的程序风格。而在不得不使用 C 时，应该注意使用 extern "C" 这种特性， 将 C 语言的代码与 C++ 代码进行分离编译，再统一链接这种做法，例如： // foo.h #ifdef __cplusplus extern "C" { #endif int add(int x, int y); #ifdef 也具有了和其他现代编程语言一样，某种意义上提供了无需操心变量类型的使用习惯。 auto auto 在很早以前就已经进入了 C++，但是他始终作为一个存储类型的指示符存在，与 register 并 存。在传统 C++ 中，如果一个变量没有声明为 register 变量，将自动被视为一个 auto 变量。而随着 18 2.3 类型推导 第 2 章语言可用性的强化 register 被弃用（在 C++17 C++ 中，如果程序员没有提供，编译器会默认为对象生成默认构造函数、复制构造、赋值 算符以及析构函数。另外，C++ 也为所有类定义了诸如 new delete 这样的运算符。当程序员有需要时， 可以重载这部分函数。 这就引发了一些需求：无法精确控制默认函数的生成行为。例如禁止类的拷贝时，必须将复制构造 函数与赋值算符声明为 private。尝试使用这些未定义的函数将导致编译或链接错误，则是一种非常不

“一本通俗易懂的資料結構與演算法入門書，引導讀者手腦並用地學習，強烈推薦演算法初學者閱讀！” ——鄧俊輝，清華大學計算機系教授 “如果我當年學資料結構與演算法時有《Hello 演算法》，學起來應該會簡單 10 倍！” ——李沐，亞馬遜資深首席科學家 電腦的出現為世界帶來了巨大的變革，它憑藉高速的運算能力與卓越的可程式化特性，成為執行演算法 與處理資料的理想媒介。無論是電玩遊戲的逼真畫面、自動駕駛的智慧決策，還是 案。從資料結構與演算法的角度看，這種方法本質上是“貪婪”演算法。 小到烹飪一道菜，大到星際航行，幾乎所有問題的解決都離不開演算法。計算機的出現使得我們能夠透過程 式設計將資料結構儲存在記憶體中，同時編寫程式碼呼叫 CPU 和 GPU 執行演算法。這樣一來，我們就能把 生活中的問題轉移到計算機上，以更高效的方式解決各種複雜問題。 Tip 如果你對資料結構、演算法、陣列和二分搜尋等概念仍感到一知半解，請繼續往下閱讀，本書將引導 各步驟都有確定的含義，在相同的輸入和執行條件下，輸出始終相同。 1.2.2 資料結構定義 資料結構（data structure）是組織和儲存資料的方式，涵蓋資料內容、資料之間關係和資料操作方法，它具 有以下設計目標。 第 1 章 初識演算法 www.hello‑algo.com 14 ‧ 空間佔用儘量少，以節省計算機記憶體。 ‧ 資料操作儘可能快速，涵蓋資料訪問、新增、刪除、更新等。 ‧ 提供簡潔的資料表示和邏輯資訊，以便演算法高效執行。

尺寸與方向：關於映像模式（座標系統） / 688 分頁 / 693 表頭（Header）與表尾（Footer）/ 695 動態計算頁碼 / 696 列印預覽（Print Preview） / 697 本章回顧 / 698 第 13 章 多重文件與多重顯示 / 701 MDI 第２章 C++ 的重要性質 85 静态成员（变量与函数） class SavingAccount { private: char m_name[40]; // 存戶姓名 char m_addr[60]; // 存戶㆞址 double m_total; // 存款額 double m_rate; // 利率 ... }; 当我们调用： ((CDocument)mydoc) 成员变量的一种方式，注意，此刻还没有诞生任何对象实体： 下面这种情况则是产生一个对象后，透过对象来处理static 成员变量： 你得搞清楚一个观念，static 成员变量并不是因为对象的实现而才得以实现，它本来就存 在，你可以想象它是一个全域变量。因此，第一种处理方式在意义上比较不会给人错误 的印象。 只要access level 允许，任何函数（包括全域函数或成员函数，static 或non-static）都可

进阶：模板元编程与函数式编程 4.编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 由于在犄角旮旯里光线反弹的次数多，算得比其 他块的慢，而有的块却算得快。但是因为木桶原 理，最后花的时间由最慢的那个线程决定，因此 变成 1 分 30 秒了，多出来的 30 秒里 1 号和 2 号 核心在闲置着，因为任务简单已经算完了，只有 4 号核心一个人在处理额外的光线。 1 2 3 4 1 分 15 秒 1 分 30 秒 0 分 45 秒 0 分 30 秒 解决 1 ：线程数量超过 CPU 核心数量，让系统调度保证各个核心始终饱和 部大区域则以类似 Z 字型的曲线遍历，这样 能保证每次访问的数据在地址上比较靠近，并 且都是最近访问过的，从而已经在缓存里可以 直接读写，避免了从主内存读写的超高延迟。 • 下次课会进一步深入探讨访存优化，详细剖析 这个案例，那么下周六 14 点敬请期待。 第 6 章：并发容器 std::vector 扩容时会移动元素 • std::vector 内部存储了一个指针，指向一段容量 capacity