深入浅出访存优化 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

）“ 2021 年 CCF 科 学技术奖科技进步卓越奖”。 伴随市场对于知识图谱应用的不断深入，图数据规模和应用性能之间的矛盾愈 加凸显，海致针对以上背景展开了系统性的技术攻关，解决了图数据的高效存 储、索引及复制难题，提出了基于图缩减的高效分析方法，并孵化出了一个大 规模图数据分析平台 AtlasGraph 。 5 获得 2022 年中国电子学会科学技术奖科技进步一等奖 中国电子学会发布的《 其中异质图建模与表示学习技术和超大规模图学习系统处于国际领 先水平。” 以终为始，以行为知，这一项目从图计算所面临的挑战出发，解决了大规模图数据所产生 的建模能力不足、结构知识难用、巨量数据难算等技术挑战，实现了大规模复杂异质图数 据的表示学习模型、语义推荐和风险管理关键技术，构建了完整的兼具理论指导与应用检 验的大规模图数据智能分析系统与平台，满足了大数据时代从复杂异质图数据中进行知识 余种图计算算法 ，可扩展的分析引擎支持更复 杂的数据挖掘和机器学习场景 MPP Massively Parallel Processing 架构，大规模集群 分布式存储及并行计 算， Shared Nothing 模式支 持存储计算分离 高性能 基于 Rust 开发的分布式存储引 擎及图计算引擎，精细的内存 管理设计，内置索引系统，支 持毫秒级的并发查询响应速度 易用 AQL(Atlas Graph

TPS 的流量 可演化性：业务逻辑与底层 API 解耦，当业务发生改变 时，底层 API 不用改变 分布式账务系统 设计理念 - Rust 是我们可靠的基石 分布式账务系统 存算分离 API 解耦 读写分离 层级账号 Rust ● 事务层与账户层分 离 ● 独立水平扩展 ● CQRS ● Event Sourcing ● 针对读场景，写场 景分别优化 ● 稳定的底层 API

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

进阶：模板元编程与函数式编程 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

进阶：模板元编程与函数式编程 4.编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 non-constexpr 函数。而且 constexpr 函数必须是内联 （ inline ）的，不能分离声明和定义在另一个文件里。标准库的很多函数如 std::min 也是 constexpr 函数，可以放心大胆在模板尖括号内使用。 模板的难题：移到另一个文件中定义 • 如果我们试着像传统函数那样分离模板函数的声明与实现： • 就会出现 undefined reference 错误： 模板的难题：移到另一个文件中定义（续） 里只看到 sumto<> 函数的两份声明，从而出错。 • 解决：在看得见 sumto<> 定义的 sumto.cpp 里，增加两个显式编译模板的声明： 一般来说，我会建议模板不要 分离声明和定义，直接写在头 文件里即可。如果分离还要罗 列出所有模板参数的排列组合 ，违背了开 - 闭原则。 模板的惰性：延迟编译 • 要证明模板的惰性，只需看这个例子： • 要是编译器哪怕细看了一眼：字符串怎么可能被写入呢？肯定是会出错的。