如需总的线程编号： blockDim * blockIdx + threadIdx 分离 __device__ 函数的声明和定义：出错 • 默认情况下 GPU 函数必须定义在同一个文件里。 如果你试图分离声明和定义，调用另一个文件里 的 __device__ 或 __global__ 函数，就会出错 。 分离 __device__ 函数的声明和定义：解决 • 开启 CMAKE_CUDA_ CMAKE_CUDA_SEPARABLE_COMPILATION 选 项（设为 ON ），即可启用分离声明和定义的支持。 • 不过我还是建议把要相互调用的 __device__ 函数放在 同一个文件，这样方便编译器自动内联优化（第四课讲 过）。 两种开启方式：全局有效 or 仅针对单个程序 只对 main 这个程序启用： 对下方所有的程序启用（推荐）： 顺便一提， CXX_STANDARD 和 CUDA_ARCHITECTURES cudaMallocManaged 、 cudaFree • 如果我没记错的话，统一内存是从 Pascal 架构开始支持的，也就是 GTX9 开头及以上 。 • 虽然方便，但并非完全没有开销，有条件的话还是尽量用分离的设备内存和主机内存吧。 第 3 章：数组 分配数组 • 如 malloc 一样，可以用 cudaMalloc 配 合 n * sizeof(int) ，分配一个大小为 n 的 整型数组。这样就会有

lock , Rows level lock , 读写性能都非常优秀 读写性能都非常优秀 • 能够承载大数据量的存储和访问 能够承载大数据量的存储和访问 • 拥有自己独立的缓冲池，能够缓存数据和索引 拥有自己独立的缓冲池，能够缓存数据和索引 MySQL 架构设计—应用架构 强一致性 对读一致性的权衡，如果是对读写实时性要求非常高的话， 就将读写都放在 M1 上面， M2 只是作为 standby 设计合理架构，如果 MySQL MySQL 访问频繁，考虑 访问频繁，考虑 Master/Slave Master/Slave 读写分离；数据库分表、数据库切片（分 读写分离；数据库分表、数据库切片（分 布式），也考虑使用相应缓存服务帮助 布式），也考虑使用相应缓存服务帮助 MySQL MySQL 缓解访问 缓解访问 压力 压力 系统优化 系统优化  配置合理的 配置合理的 MySQL 1024 MySQL 服务器同时处理的数据库连接的最大 数量 query_cache_size 0 ( 不打开 ） 128M 查询缓存区的最大长度，按照当前需求，一 倍一倍增加，本选项比较重要 sort_buffer_size 512K 128M 每个线程的排序缓存大小，一般按照内存可 以设置为 2M 以上，推荐是 16M ，该选项对 排序 order by ， group by 起作用 record_buffer

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 通常认为利用同时处理 4 个 float 的 SIMD 指令可以加速 4 倍。但是如果你的算法不 适合 SIMD ，则可能加速达不到 4 倍；也有因为 SIMD 让访问内存更有规律，节约了指 令解码和指令缓存的压力等原因，出现加速超过 4 倍的情况。 第 1 章：化简 编译器优化：代数化简 编译器优化：常量折叠 编译器优化：举个例子 编译器优化：我毕竟不是万能的 结论：尽量避免代码复杂化，避免使用会造 而我们可以用 const 禁止写入访问。 结论：所有非 const 的指针都声明 __restrict 。 禁止优化： volatile 结论：加了 volatile 的对象，编 译器会放弃优化对他的读写操作 。 做性能实验的时候非常有用。 注意一下区别 1. volatile int *a 或 int volatile *a 2. int *__restrict a • 语法上区别： volatile

com/parallel101/course 为什么往 int 数组里赋值 1 比赋值 0 慢一倍？ 第 1 章：内存带宽 cpu-bound 与 memory-bound • 通常来说，并行只能加速计算的部分，不能加速内存读写的部分 。 • 因此，对 fill 这种没有任何计算量，纯粹只有访存的循环体，并 行没有加速效果。称为内存瓶颈（ memory-bound ）。 • 而 sine 这种内部需要泰勒展开来计算，每次迭代计算量很大的 在太少了。 • 计算太简单，数据量又大，并行只带来了多线程调度的额外开销 。 • 小彭老师经验公式： 1 次浮点读写 ≈ 8 次浮点加法 • 如果矢量化成功（ SSE ）： 1 次浮点读写 ≈ 32 次浮点加法 • 如果 CPU 有 4 核且矢量化成功： 1 次浮点读写 ≈ 128 次浮点加 法 常见操作所花费的时间 • 图中加法 (add) 和乘法 (mul) 都指的整数。 • • 区别是浮点的乘法和加法基本是一样速度。 • L1/2/3 read 和 Main RAM read 的时间指的是 读一个缓存行（ 64 字节）所花费的时间。 • 根据计算： 125/64*4≈8 • 即从主内存读取一次 float 花费 8 个 cycle ， 符合小彭老师的经验公式。 • “right” 和“ wrong” 指的是分支预测是否成功。 多少计算量才算多？ • 看右边的

疏网格、位运算、浮点的二进制格式、内存带宽优 化 面向人群：图形学、 CFD 仿真、深度学习编程人 员 第 0 章：稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 来存储 分离 read/write/create 三种访问模式 foreach 直接给出当前坐标指向的值 改用 unordered_map 来存储 unordered_map 手动 read(i, j) 也一样速度 更高效。其实 sizeof(std::mutex) = 40 字节，而 sizeof(tbb::spin_mutex) = 1 字节…… 小彭老师解决：访问者模式 把写入过的块地址缓存起来，可以避免多次访问全局表的开销。缓存在访问 者 (accessor) 的成员 map 里。访问者对象被我用 OpenMP 标记为 firstprivate ，意味着这个 map 是线程局部的，因此对他的访问不需要加锁， parallel for collapse(2) 遍历二维区间。 把 func 捕获为 firstprivate ，从而支持用 lambda 捕获的访问者模式。 实现访问者模式 • 额，总之就是每一层都有一个缓存。 第 5 章：量化整型 使用 int ：每个占据 4 字节 • 记得我第七课说过，一个简单的循环体往 往会导致内存成为瓶颈（ memory- bound ）。 • 右边就是一个很好的例子。

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 。 • 所以， download 函数才会出师未捷身先死 ——还没开始执行他的线程就被销毁了。 解构函数不再销毁线程： t1.detach() • 解决方案：调用成员函数 detach() 分离该 线程——意味着线程的生命周期不再由当 前 std::thread 对象管理，而是在线程退 出以后自动销毁自己。 • 不过这样还是会在进程退出时候自动退出 。 解构函数不再销毁线程：移动到全局线程池 上仍是 const 的。因此，为了让 this 为 const 时仅仅给 m_mtx 开后门，可以用 mutable 关键字修饰他，从而所有成员里 只有他不是 const 的。 为什么需要读写锁？ • 刚才说过 mutex 就像厕所，同一时刻只有一个人能上。但是如果“上”有两种方式呢？ • 假设在平行世界，厕所不一定是用来拉的，还可能是用来喝的（只是打个比方，请勿尝试） • 喝厕所里的水时，可以多个人插着吸管一起喝。

，可扩展的分析引擎支持更复 杂的数据挖掘和机器学习场景 MPP Massively Parallel Processing 架构，大规模集群 分布式存储及并行计 算， Shared Nothing 模式支 持存储计算分离 高性能 基于 Rust 开发的分布式存储引 擎及图计算引擎，精细的内存 管理设计，内置索引系统，支 持毫秒级的并发查询响应速度 易用 AQL(Atlas Graph Query Language) 实际执行时，执行器等待流数据，处 理后将数据推送到下一个执行器 切分执行计划，将执行计划划分成不 同的执行阶段 内存缓存结构：加速图数据查询 • 由于图数据的查询通常是 IO 密集型，且访问的数据随机又分散，拥有内存缓存能起到很 好的加速效果 • 要想让内存缓存发挥最大的作用，就要能在有限的内存中存下尽量多的图数据 • 例如，对于属性的存储，可以通过自行序列化 / 反序列化大幅节省内存 u32 u32 string string 定长 变长 高可用技术方案 基于 Chain Replication （ CRAQ ） 算法实现，进行数据副本处理，头 结点写，多结点读，支持读写分离 ，提供更好的并发查询能力 数据高可用实现 Chain Replication 数据高可用方案 服务高可用实现 系统中 Meta ， TS 服务采用主备架 构，基于 Raft 算法实现租约，进行

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 这种情况出现时，就意味着你需要把成员变量的读写封装为成员函数 不变性：请勿滥用封装 • 仅当出现“修改一个成员时，其他也成员要 被修改，否则出错”的现象时，才需要 getter/setter 封装。 • 各个成员之间相互正交，比如数学矢量类 Vec3 ，就没必要去搞封装，只会让程序员 变得痛苦，同时还有一定性能损失：特别 是如果 getter/setter 函数分离了声明和定 义，实现在另一个文件时！ 移动进阶：交换两者的值 • 除了 std::move 可以把 v2 移动到 v1 外 ， • 还可以通过 std::swap 交换 v1 和 v2 。 • swap 在高性能计算中可以用来实现双缓存 （ ping-pong buffer ）。 swap 可能是 这样实现的 ： 还有哪些情况会触发“移动” • 这些情况下编译器会调用移动： • return v2

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 显然不是。甚至在两个处理器上同时运行两个线程也不见得可以获得两倍的性能。相似的 ，大多数多线程的应用不会比双核处理器的两倍快。他们应该比单核处理器运行的快，但 是性能毕竟不是线性增长。 • 为什么无法做到呢？首先，为了保证缓存一致性以及其他握手协议需要运行时间开销。在 今天，双核或者四核机器在多线程应用方面，其性能不见得的是单核机器的两倍或者四倍。 这一问题一直伴随 CPU 发展至今。 并发和并行的区别 • 运用多线程的方式和动机，一般分为两种。 ）将矩阵进行分块。块内部小区 域按照常规的两层循环访问以便矢量化，块外 部大区域则以类似 Z 字型的曲线遍历，这样 能保证每次访问的数据在地址上比较靠近，并 且都是最近访问过的，从而已经在缓存里可以 直接读写，避免了从主内存读写的超高延迟。 • 下次课会进一步深入探讨访存优化，详细剖析 这个案例，那么下周六 14 点敬请期待。 第 6 章：并发容器 std::vector 扩容时会移动元素 • std::vector

cudaSurfaceObject_t ）。 • 考虑到多维数组始终是需要通过表面对象来访问的，这 里我们让表面对象继承自多维数组。 • 在核函数中可以用 surf3Dread 和 surf3Dwrite 来读写 表面对象中的元素， x,y,z 参数指定要访问元素的坐标 ，要注意 x 必须乘以 sizeof( 元素类型 ) ，否则出错。 • 这里用了访问者模式（ Accessor ， GPU 编程常用）。 编程常用）。 原来的 CudaSurface 管理资源，禁止拷贝。然后单独 弄一个访问者类 CudaSurfaceAccessor ，不管理资源 ，仅仅是指向资源的一个弱引用，可以随意拷贝。并把 读写访问的方法（ surf3Dread ）定义在访问者类。 CUDA 表面对象：封装 • 此外，表面对象还支持自动判断 x,y,z 坐标是否越界 ， surf3Dread/write 的最后一个参数，用于指定出现 新采样。核函数的 gridDim 通过上整除技巧保证每个元素都能访问到， blockDim 为 8x8x8=512 。 • 如果在 resample_kernel 需要读取 clr ，然后再写入 clr ，并且读写是不同的坐标位置。 • 因此对 clr 和 vel 使用了双缓冲，写入 clrNext 的同时读取 clr 没有冲突，写入完毕后对调 clrNext 和 clr 。 投影部分 投影部分 •