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

而我们可以用 const 禁止写入访问。 结论：所有非 const 的指针都声明 __restrict 。 禁止优化： volatile 结论：加了 volatile 的对象，编 译器会放弃优化对他的读写操作 。 做性能实验的时候非常有用。 注意一下区别 1. volatile int *a 或 int volatile *a 2. int *__restrict a • 语法上区别： volatile a 总是对齐到 16 字节，在 GCC 编译器中这样 写： 但这样不通用，因此 C++20 引入了标准化的 std::assume_aligned ： movups 变成了 movaps 对齐的读写可能 带来微乎其微的 性能提升…… 数组求和： reduction 的优化 你看懂了吗？没关系！小彭老师也没看 懂！总之非常高效就对了！ 第 5 章：循环 循环中的矢量化：还在伺候指针别名 结构体的内存布局： AOS 与 SOA • AOS （ Array of Struct ）单个对象的属性紧挨着存 • xyzxyzxyzxyz • SOA （ Struct of Array ）属性分离存储在多个数组 • xxxxyyyyzzzz • AOS 必须对齐到 2 的幂才高效， SOA 就不需要。 • AOS 符合直觉，不一定要存储在数组这种线性结构， 而 SOA 可能无法保证多个数组大小一致。

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) 都指的整数。 • 用了 6 核才饱和。 • 结论：要想利用全部 CPU 核心，避免 mem-bound ，需要 func 里有足够的计算 量。 • 当核心数量越多， CPU 计算能力越强，相 对之下来不及从内存读写数据，从而越容 易 mem-bound 。 1 2 4 6 8 10 0 50 100 150 200 250 300 350 funcA funcB funcC 内存信息查看工具：

。 • 所以， download 函数才会出师未捷身先死 ——还没开始执行他的线程就被销毁了。 解构函数不再销毁线程： t1.detach() • 解决方案：调用成员函数 detach() 分离该 线程——意味着线程的生命周期不再由当 前 std::thread 对象管理，而是在线程退 出以后自动销毁自己。 • 不过这样还是会在进程退出时候自动退出 。 解构函数不再销毁线程：移动到全局线程池 上仍是 const 的。因此，为了让 this 为 const 时仅仅给 m_mtx 开后门，可以用 mutable 关键字修饰他，从而所有成员里 只有他不是 const 的。 为什么需要读写锁？ • 刚才说过 mutex 就像厕所，同一时刻只有一个人能上。但是如果“上”有两种方式呢？ • 假设在平行世界，厕所不一定是用来拉的，还可能是用来喝的（只是打个比方，请勿尝试） • 喝厕所里的水时，可以多个人插着吸管一起喝。 结论：读可以共享，写必须独占，且写和读不能共存。 • 针对这种更具体的情况，又发明了读写锁，他允许的状态有： 1. n 个人读取，没有人写入。 2. 1 个人写入，没有人读取。 3. 没有人读取，也没有人写入。 读写锁： shared_mutex • 为此，标准库提供了 std::shared_mutex 。 • 上锁时，要指定你的需求是拉还是喝，负 责调度的读写锁会帮你判断要不要等待。 • 这里 push_back()

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 。 投影部分 投影部分 •

疏网格、位运算、浮点的二进制格式、内存带宽优 化 面向人群：图形学、 CFD 仿真、深度学习编程人 员 第 0 章：稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 来存储 分离 read/write/create 三种访问模式 foreach 直接给出当前坐标指向的值 改用 unordered_map 来存储 unordered_map 手动 read(i, j) 也一样速度 int64_t ：每个占据 8 字节 • 如果用更大的数据类型，用时会直接提升两倍！ • 这是因为 i % 2 的计算时间，完全隐藏在内存 的超高延迟里了。 • 可见，当数据量足够大，计算量却不多时，读写 数据量的大小唯一决定着你的性能。 • 特别是并行以后，计算量可以被并行加速，而访 存却不行。 使用 int8_t ：每个占据 1 字节 • 因此我们可以把数据类型变小，这样所需的内存 量就变小，从而内存带宽也可以减小！ 类操作系统的 mmap 函数，他位于 sys/mman.h 头文件里。 • Windows 可以用 VirtualAllocateEx 之类。 • mmap 出来的起始地址保证是对齐到 4KB 的，读写访问其 中偏移地址时，会按页的粒度自动分配和释放内存，从而满 足稀疏数据结构“按需分配”的需求。且由于分页是硬件自动 来做的，比我们软件哈希和指针数组的稀疏更高效，写起来 就和普通的二维数组没什么两样，就好像顺序访问。也用不

之间的区别）。在编写 C++ 时，也应该尽可能 的避免使用诸如 void* 之类的程序风格。而在不得不使用 C 时，应该注意使用 extern "C" 这种特性， 将 C 语言的代码与 C++ 代码进行分离编译，再统一链接这种做法，例如： // foo.h #ifdef __cplusplus extern "C" { #endif int add(int x, int y); #ifdef 和 flag 而言，他们在两个并行的线程中被读写，出现了竞争。除此之外，即便我们忽略竞争 读写，仍然可能受 CPU 的乱序执行，编译器对指令的重排的影响，导致 a = 5 发生在 flag = 1 之后。 从而 b 可能输出 0。 68 7.5 原子操作与内存模型 第 7 章并行与并发 原子操作 std::mutex 可以解决上面出现的并发读写的问题，但互斥锁是操作系统级的功能，这是因为一个互 关于同步条件的研究有着非常久远的历史，我们在这里不进行赘述。读者应该明白，现代 CPU 体系 结构提供了 CPU 指令级的原子操作，因此在 C++11 中多线程下共享变量的读写这一问题上，还引入了 std::atomic 模板，使得我们实例化一个原子类型，将一个原子类型读写操作从一组指令，最小化到单 个 CPU 指令。例如： std::atomic counter; 并为整数或浮点数的原子类型提供了基本的数值成员函数，举例来说，包括

去，并把DstDir 中多出来的文件删除，使ScrDir 和DstDir 的文件保 33 持完全相同。之所以不做xcopy 完全拷贝动作，为的是节省拷贝时间（做为备份装置， 通常是软盘或磁带或可擦写光盘MO，读写速度并不快）。 JBACKUP 没有能力处理SrcDir 底下的子目录文件。如果要处理子目录，漂亮的作法是 使用递归（recursive），但是有点伤脑筋，这一部份留给你了。我的打字速度还算快，多 对象导向术语中有一个名为persistence，意思是永续存留。放在RAM 中的东西，生命 受到电力的左右，不可能永续存留；唯一的办法是把它写到文件去。MFC 的一个术语 Serialize，就是做有关文件读写的永续存留动作，并且实做作出一个虚拟函数，就叫作 Serialize。 看起来永续存留与本节的主题「动态生成」似乎没有什么干连。有！你把你的资料储存 到文件，这些资料很可能（通常是）对象中的成员变量 有一套Serialize 机制，目的在于把档名的选择、文件的开关、缓冲区的建立、资 料的读写、萃取运算子（>>）和嵌入运算子（<<）的多载（overload）、对象的动态生成 ... 都包装起来。 上述Serialize 的各部份工作，除了资料的读写和对象的动态生成，其余都是支节。动态 生成的技术已经解决，让我们集中火力，分析资料的读写动作。 第３章 MFC 六大關鍵技術之模擬 161 CObject

这种情况出现时，就意味着你需要把成员变量的读写封装为成员函数 不变性：请勿滥用封装 • 仅当出现“修改一个成员时，其他也成员要 被修改，否则出错”的现象时，才需要 getter/setter 封装。 • 各个成员之间相互正交，比如数学矢量类 Vec3 ，就没必要去搞封装，只会让程序员 变得痛苦，同时还有一定性能损失：特别 是如果 getter/setter 函数分离了声明和定 义，实现在另一个文件时！

non-constexpr 函数。而且 constexpr 函数必须是内联 （ inline ）的，不能分离声明和定义在另一个文件里。标准库的很多函数如 std::min 也是 constexpr 函数，可以放心大胆在模板尖括号内使用。 模板的难题：移到另一个文件中定义 • 如果我们试着像传统函数那样分离模板函数的声明与实现： • 就会出现 undefined reference 错误： 模板的难题：移到另一个文件中定义（续） 里只看到 sumto<> 函数的两份声明，从而出错。 • 解决：在看得见 sumto<> 定义的 sumto.cpp 里，增加两个显式编译模板的声明： 一般来说，我会建议模板不要 分离声明和定义，直接写在头 文件里即可。如果分离还要罗 列出所有模板参数的排列组合 ，违背了开 - 闭原则。 模板的惰性：延迟编译 • 要证明模板的惰性，只需看这个例子： • 要是编译器哪怕细看了一眼：字符串怎么可能被写入呢？肯定是会出错的。 1. 当函数有多条 return 语句时，所有语句的返回类型必须一致，否则 auto 会报错。 2. 当函数没有 return 语句时， auto 会被推导为 void 。 3. 如果声明和实现分离了，则不能声明为 auto 。比如： auto func(); // 错误 C++ 特性：引用（ int & ） • 众所周知， C++ 中有一种特殊的类型，叫做引用。只需要在原类型后面加一个 &