上线资源设备只能人工确认，资源设备没有统 一的管理平台。 研发排查问题困难，对于服务进行诊断，每有 一个资源设备都需要给研发单独分配 SSH 权 限，管理成本巨高。 开发 debug 过程需要登录统一的内网主机使 用 Kubectl 操作，权限不可控，风险大。 对于新上项目，面对不同的使用场景，需要创建多 条 Jenkins Job ，配置繁琐，维护负担重。 与 传 统 的 业 务 研

内部引入了一片极小的存储 器——虽然小，但是读写速度却特别快。这片小而快的 存储器称为缓存（ cache ）。 • 当 CPU 访问某个地址时，会先查找缓存中是否有对应的 数据。如果没有，则从内存中读取，并存储到缓存中； 如果有，则直接使用缓存中的数据。 • 这样一来，访问的数据量比较小时，就可以自动预先加 载到这个更高效的缓存里，然后再开始做运算，从而避 免从外部内存读写的超高延迟。 缓存的分级结构 回缓存中的数据。如果找不到，则向主内存发送请求，等读 取到该地址的数据，就创建一个新条目。 • 在 x86 架构中每个条目的存储 64 字节的数据，这个条目 又称之为缓存行（ cacheline ）。 • 当访问 0x0048~0x0050 这 4 个字节时，实际会导致 0x0040~0x0080 的 64 字节数据整个被读取到缓存中。 • 这就是为什么我们喜欢把数据结构的起始地址和大小对齐到 64 如有多级缓存，则一级缓存失效后会丢给二级缓存。 连续访问与跨步访问 • 如果访问数组时，按一定的间距跨步访问，则效率如何？ • 从 1 到 16 都是一样快的， 32 开始才按 2 的倍率变慢，为什么？ • 因为 CPU 和内存之间隔着缓存，而缓存和内存之间传输数据的最小 单位是缓存行（ 64 字节）。 16 个 float 是 64 字节，所以小于 64 字节的跨步访问，都会导致数据全部被读取出来。而超过

线程，用于处理大吞吐量的数据。 获取线程编号 • 可以通过 threadIdx.x 获取当前线程的编 号，我们打印一下试试看。 • 这是 CUDA 中的特殊变量之一，只有在 核函数里才可以访问。 • 可以看到线程编号从 0 开始计数，打印出 了 0 ， 1 ， 2 。这也是我们指定了线程数 量为 3 的缘故。 • 等等，为什么后面有个 .x ？稍后再说明。 获取线程数量 • 还可以用 int 。 • 可以通过 cudaGetErrorName 获取该 enum 的具体名 字。这里显示错误号为 77 ，具体名字是 cudaErrorIllegalAddress 。意思是我们访问了非法的地 址，和 CPU 上的 Segmentation Fault 差不多。 封装好了： helper_cuda.h • 其实 CUDA toolkit 安装时，会默认附带一系列案例代码， 报告出错所在的行号，函数名等，很方便。 堆上分配试试？ • 那你可能会想，难道是因为我的 ret 创建 在栈上，所以 GPU 不能访问，才出错的 ？ • 于是你试图用 malloc 在堆上分配一个 int 来给 GPU 访问，结果还是失败了。 原因： GPU 使用独立的显存，不能访问 CPU 内存 • 原来， GPU 和 CPU 各自使用着独立的内 存。 CPU 的内存称为主机内存 (host)

面向人群：图形学、 CFD 仿真、深度学习编程人 员 第 0 章：稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 来存储 分离 read/write/create 三种访问模式 foreach 直接给出当前坐标指向的值 改用 unordered_map 来存储 unordered_map 手动 read(i, j) 也一样速度 索性把坐标和值打包成 tuple Python 的模运算 a % b 的值始终是 [0, b) 区间内的正数，非常方便。 对稀疏数据结构造成的问题 • 如果这里的 x 是负数，则 x % B 也是负数，会造成对 m_block 的越界访问。 • 因此 % 会返回负数对 CFD 用户来说是个很大的坑点，很多人想当然地用 % 做循环边界， 然而这对负方向会不起作用。 解决： (a % b + b) % b • 我看一些 CFD 用户喜欢写 pdf 第 4 章：并行与随机访问 回到指针的数组 试图并行地访问：出错了 为什么？因为多个核心同时访问了 m_block ，造成数据竞争。所 以有的指针被重复分配了两遍，写入了那个地址却没有实际被存到 m_data 这个指针数组里。因此结果不对，还造成了内存泄露。 解决：使用互斥量和原子变量 暴力解决方案就是用 std::mutex 避免多个线程同时访问。 然而这样会严重影响性能，锁和原子多了，就根本并行不起来。

字的长度除了决定一次处理的整数大小之外，还决定了能访问的内存地址的范围。 • 这是因为内存是一维排列的，假如内存容量是 65536 字节，那所谓的内存地址实际上就 是一个从 0 到 65535 范围的整数，也就是两个字节组成的字。 • 处理器去读写内存的时候靠的是寄存器提供的地址，因此寄存器的大小（也就是字的大 小）决定了他能读写的内存大小，例如： • 由于 16 位计算机的寄存器只能存储 16 位，他只能访问 65536 字节（ ）的内存 。 • 由于 32 位计算机的寄存器只能存储 32 位，他只能访问 4 GB 的内存。 • 由于 64 位计算机的寄存器能存储 64 位，他理论上能访问 16777216 TB 的内存！ • 因此，如果你的电脑内存超过了 4 GB ，那肯定是 32 位电脑不用说了。 • 而 64 位计算机理论上能访问如此大量的内存，虽然目前看来是用不到。 知识拓展 • 虽然 64 位计算机的寄存器能处理 x64 架构实际上只能访 问 512GB 内存，如果插了超过这个大小的内存条他也不会认出来。 • 此外， 16 位计算机实际上能通过额外的段寄存器访问到 20 位的内存地址（ 1MB ）。 • 32 位计算机还能通过 PAE 技术（物理地址扩展）访问到 36 位的内存地址（ 64GB ） 。 • 64 位计算机反而是因为 16777216 TB 太大，内存地址被阉割到了 39 位（ 512GB

的三维数组之间拷贝数据。 CUDA 表面对象：封装 • 要访问一个多维数组，必须先创建一个表面对象 （ cudaSurfaceObject_t ）。 • 考虑到多维数组始终是需要通过表面对象来访问的，这 里我们让表面对象继承自多维数组。 • 在核函数中可以用 surf3Dread 和 surf3Dwrite 来读写 表面对象中的元素， x,y,z 参数指定要访问元素的坐标 ，要注意 x 必须乘以 sizeof( sizeof( 元素类型 ) ，否则出错。 • 这里用了访问者模式（ Accessor ， GPU 编程常用）。 原来的 CudaSurface 管理资源，禁止拷贝。然后单独 弄一个访问者类 CudaSurfaceAccessor ，不管理资源 ，仅仅是指向资源的一个弱引用，可以随意拷贝。并把 读写访问的方法（ surf3Dread ）定义在访问者类。 CUDA 表面对象：封装 • 此外，表面对象还支持自动判断 ；对于写来说越界会放弃写入，不修改数组中的任 何值。 • 表面对象保障了高效的访存，并且自动判断越界，体 现了 GPU 作为图形学专业硬件的能力。 CUDA 纹理对象：封装 • 表面对象访问数组是可读可写的。纹理对象也可以访问 数组，不过是只读的。好处是他可以通过浮点坐标来访 问，且提供了线性滤波的能力。 • 在核函数中可以通过 tex3D 来读取纹理中的值。 • 之所以纹理是因为 GPU 一开始是渲染图形的专用硬件

4 。 • size_t size() const noexcept; vector 容器： operator[] • 要访问 vector 里的元素，只需用 [] 运算符 ： • 例如 a[0] 访问第 0 个元素（人类的第一 个） • 例如 a[1] 访问第 1 个元素（人类的第二 个） • int &operator[](size_t i) noexcept; • int const const noexcept; vector 容器： operator[] • 值得注意的是， [] 运算符在索引超出数组大 小时并不会直接报错，这是为了性能的考虑。 • 如果你不小心用 [] 访问了越界的索引，可能 会覆盖掉别的变量导致程序行为异常，或是访 问到操作系统未映射的区域导致奔溃。 • int &operator[](size_t i) noexcept; • int const 存储的数组，因此只要得到了首地址，下一 个元素的地址只需指针 +1 即可。 • 因为指针的 p[i] 相当于 *(p + i) ，因此可以 把 data() 返回的首地址指针当一个数组来 访问。 • int *data() noexcept; • int const *data() const noexcept; vector 容器： data() 获取首地址指针 • data()

tbb::simple_partitioner 能够按照给定的粒度 大小（ grain ）将矩阵进行分块。块内部小区 域按照常规的两层循环访问以便矢量化，块外 部大区域则以类似 Z 字型的曲线遍历，这样 能保证每次访问的数据在地址上比较靠近，并 且都是最近访问过的，从而已经在缓存里可以 直接读写，避免了从主内存读写的超高延迟。 • 下次课会进一步深入探讨访存优化，详细剖析 这个案例，那么下周六 • 而 grow_by(n) 则可以一次扩充 n 个元素。 他同样是返回一个迭代器（ iterator ），之 后可以通过迭代器的 ++ 运算符依次访问 连续的 n 个元素， * 运算符访问当前指 向的元素。 可安全地被多线程并发访问 • 除了内存不连续、指针和迭代器不失效的 特点， tbb::concurrent_vector 还是一个多 线程安全的容器，能够被多个线程同时并 不建议通过索引随机访问 • 因为 tbb::concurrent_vector 内存不连续 的特点，通过索引访问，比通过迭代器访 问的效率低一些。 • 因此不推荐像 a[i] 这样通过索引随机访问 其中的元素， *(it + i) 这样需要迭代器跨步 访问的也不推荐。 推荐通过迭代器顺序访问 • 最好的方式是用 begin() 和 end() 的迭代 器区间，按顺序访问。 parallel_for

了单独一个指令。这里尽管不 是地址，但同样可以利用 lea 指令简化生成的代码大小。 eax = rdi + rsi * 8 指针访问对象：线性访问地址 rsi = (int64_t)esi eax = *(int *)(rdi + rsi * 4) 为什么乘以 4 ？因为访问的 对象， int 的大小是 4 。 指针的索引：尽量用 size_t eax = *(int *)(rdi + rsi 的代码，从而增强你程序的吞吐能力！ • 通常认为利用同时处理 4 个 float 的 SIMD 指令可以加速 4 倍。但是如果你的算法不 适合 SIMD ，则可能加速达不到 4 倍；也有因为 SIMD 让访问内存更有规律，节约了指 令解码和指令缓存的压力等原因，出现加速超过 4 倍的情况。 第 1 章：化简 编译器优化：代数化简 编译器优化：常量折叠 编译器优化：举个例子 编译器优化：我毕竟不是万能的 编译器保证：这些指针之间不会发生重叠！ 从而他可以放心地优化成功： __restrict 关键字：只需加在非 const 的即可 实际上， __restrict 只需要加在所有具有写入 访问的指针（这里是 c ）上，就可以优化成功 。 而我们可以用 const 禁止写入访问。 结论：所有非 const 的指针都声明 __restrict 。 禁止优化： volatile 结论：加了 volatile 的对象，编 译器会放弃优化对他的读写操作

少碎片、支持大文件、能够进行索引压缩 少碎片、支持大文件、能够进行索引压缩 • 二进制层次的文件可以移植 二进制层次的文件可以移植 (Linux (Linux   Windows) Windows) • 访问速度飞快，是所有 访问速度飞快，是所有 MySQL MySQL 文件引擎中速度最快的 文件引擎中速度最快的 • 不支持一些数据库特性，比如 事务、外键约束等 不支持一些数据库特性，比如 事务、外键约束等 • 事务、外键约束等数据库特性 • Rows level lock , Rows level lock , 读写性能都非常优秀 读写性能都非常优秀 • 能够承载大数据量的存储和访问 能够承载大数据量的存储和访问 • 拥有自己独立的缓冲池，能够缓存数据和索引 拥有自己独立的缓冲池，能够缓存数据和索引 MySQL 架构设计—应用架构 强一致性 对读一致性的权衡，如果是对读写实时性要求非常高的话， 设计合理架构，如果 设计合理架构，如果 MySQL MySQL 访问频繁，考虑 访问频繁，考虑 Master/Slave Master/Slave 读写分离；数据库分表、数据库切片（分 读写分离；数据库分表、数据库切片（分 布式），也考虑使用相应缓存服务帮助 布式），也考虑使用相应缓存服务帮助 MySQL MySQL 缓解访问 缓解访问 压力 压力 系统优化 系统优化  配置合理的 配置合理的