如需总的线程编号： 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 的 整型数组。这样就会有

编译器如何自动优化：从汇编角度看 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 就像厕所，同一时刻只有一个人能上。但是如果“上”有两种方式呢？ • 假设在平行世界，厕所不一定是用来拉的，还可能是用来喝的（只是打个比方，请勿尝试） • 喝厕所里的水时，可以多个人插着吸管一起喝。

编译器如何自动优化：从汇编角度看 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

去，并把DstDir 中多出来的文件删除，使ScrDir 和DstDir 的文件保 33 持完全相同。之所以不做xcopy 完全拷贝动作，为的是节省拷贝时间（做为备份装置， 通常是软盘或磁带或可擦写光盘MO，读写速度并不快）。 JBACKUP 没有能力处理SrcDir 底下的子目录文件。如果要处理子目录，漂亮的作法是 使用递归（recursive），但是有点伤脑筋，这一部份留给你了。我的打字速度还算快，多 配置「执行线程对象」，其handle 将成为CreateThread 的传回值。 2. 设定计数值为1。 3. 配置执行线程的context。 4. 保留执行线程的堆栈。 5. 将context 中的堆栈指针缓存器（SS）和指令指针缓存器（IP）设定妥当。 看看上面的态势，的确可以显示出执行线程是CPU 分配时间的单位。所谓工作切换（context switch）其实就是对执行线程的context 的切换。 程序 对象导向术语中有一个名为persistence，意思是永续存留。放在RAM 中的东西，生命 受到电力的左右，不可能永续存留；唯一的办法是把它写到文件去。MFC 的一个术语 Serialize，就是做有关文件读写的永续存留动作，并且实做作出一个虚拟函数，就叫作 Serialize。 看起来永续存留与本节的主题「动态生成」似乎没有什么干连。有！你把你的资料储存 到文件，这些资料很可能（通常是）对象中的成员变量

编译器如何自动优化：从汇编角度看 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 。 投影部分 投影部分 •

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