1/66Bridging the Gap: Writing Portable Programs for CPU and GPU using CUDA Thomas Mejstrik Sebastian Woblistin 2/66Content 1 Motivation Audience etc.. Cuda crash course Quiz time 2 Patterns Oldschool afterwards7/66 Motivation Patterns The dark path Cuda proposal Thank you Why write programs for CPU and GPU Difference CPU/GPU Algorithms are designed differently Latency/Throughput Memory bandwidth Number of talk7/66 Motivation Patterns The dark path Cuda proposal Thank you Why write programs for CPU and GPU Difference CPU/GPU Why it makes sense? Library/Framework developers Embarrassingly parallel algorithms

1. 关于本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 1 0.2. 如何使用本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 3 0.3. 小结 . . . . . 会与不会之间反复横跳，那么这本书就是为你而写！ 如果您是「算法老手」，已经积累一定刷题量，接触过大多数题型，那么本书可以帮助你回顾与梳理算法知识 体系，仓库源代码可以被当作“刷题工具库”或“算法字典”来使用。 如果您是「算法大佬」，希望可以得到你的宝贵意见建议，或者一起参与创作。 � 前置条件 您需要至少具备任一语言的编程基础，能够阅读和编写简单代码。 0.1.2. 内容结构 本书主要内容有： 组、链表、栈、队列、 散列表、树、堆、图等数据结构，内容包括定义、优劣势、常用操作、常见类型、典型应用、实现方法等。 ‧ 算法：查找算法、排序算法、搜索与回溯、动态规划、分治算法，内容包括定义、使用场景、优劣势、时 空效率、实现方法、示例题目等。 0. 写在前面 hello‑algo.com 2 Figure 0‑1. Hello 算法内容结构 0.1.3. 致谢 本书的成书过程中，我获得了许多人的帮助，包括但不限于：

1. 关于本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 1 0.2. 如何使用本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 3 0.3. 小结 . . . . . 会与不会之间反复横跳，那么这本书就是为你而写！ 如果您是「算法老手」，已经积累一定刷题量，接触过大多数题型，那么本书可以帮助你回顾与梳理算法知识 体系，仓库源代码可以被当作“刷题工具库”或“算法字典”来使用。 如果您是「算法大佬」，希望可以得到你的宝贵意见建议，或者一起参与创作。 � 前置条件 您需要至少具备任一语言的编程基础，能够阅读和编写简单代码。 0.1.2. 内容结构 本书主要内容有： 组、链表、栈、队列、 散列表、树、堆、图等数据结构，内容包括定义、优劣势、常用操作、常见类型、典型应用、实现方法等。 ‧ 算法：查找算法、排序算法、搜索与回溯、动态规划、分治算法，内容包括定义、使用场景、优劣势、时 空效率、实现方法、示例题目等。 0. 写在前面 hello‑algo.com 2 Figure 0‑1. Hello 算法内容结构 0.1.3. 致谢 本书的成书过程中，我获得了许多人的帮助，包括但不限于：

章：内存带宽 cpu-bound 与 memory-bound • 通常来说，并行只能加速计算的部分，不能加速内存读写的部分 。 • 因此，对 fill 这种没有任何计算量，纯粹只有访存的循环体，并 行没有加速效果。称为内存瓶颈（ memory-bound ）。 • 而 sine 这种内部需要泰勒展开来计算，每次迭代计算量很大的 循环体，并行才有较好的加速效果。称为计算瓶颈（ cpu- bound 计算太简单，数据量又大，并行只带来了多线程调度的额外开销 。 • 小彭老师经验公式： 1 次浮点读写 ≈ 8 次浮点加法 • 如果矢量化成功（ SSE ）： 1 次浮点读写 ≈ 32 次浮点加法 • 如果 CPU 有 4 核且矢量化成功： 1 次浮点读写 ≈ 128 次浮点加 法 常见操作所花费的时间 • 图中加法 (add) 和乘法 (mul) 都指的整数。 • 区别是浮点的乘法和加法基本是一样速度。 • funcA 用了 2 核就饱和。 • funcB 用了 4 核才饱和。 • funcC 用了 6 核才饱和。 • 结论：要想利用全部 CPU 核心，避免 mem-bound ，需要 func 里有足够的计算 量。 • 当核心数量越多， CPU 计算能力越强，相 对之下来不及从内存读写数据，从而越容 易 mem-bound 。 1 2 4 6 8 10 0 50 100 150

• 这是因为 GPU 和 CPU 之间的通信，为了高效，是异 步的。也就是 CPU 调用 kernel<<<1, 1>>>() 后，并不 会立即在 GPU 上执行完毕，再返回。实际上只是把 kernel 这个任务推送到 GPU 的执行队列上，然后立即 返回，并不会等待执行完毕。 • 因此可以调用 cudaDeviceSynchronize() ，让 CPU 陷 入等待，等 GPU 完成队列的所有任务后再返回。从而 完成队列的所有任务后再返回。从而 能够在 main 退出前等到 kernel 在 GPU 上执行完。 定义在 GPU 上的设备函数 • __global__ 用于定义核函数，他在 GPU 上执行，从 CPU 端通过三重尖括号语法调 用，可以有参数，不可以有返回值。 • 而 __device__ 则用于定义设备函数，他在 GPU 上执行，但是从 GPU 上调用的，而 且不需要三重尖括号，和普通函数用起来一 符号，和性能优化意义上的内联无关。 • 优化意义上的内联指把函数体直接放到调用者那里去。 • 因此 CUDA 编译器提供了一个“私货”关键字： __inline__ 来 声明一个函数为内联。不论是 CPU 函数还是 GPU 都可以使 用，只要你用的 CUDA 编译器。 GCC 编译器相应的私货则 是 __attribute__((“inline”)) 。 • 注意声明为 __inline__ 不一定就保证内联了，如果函数太大编

2023 年 5 月 7 日- ff6ee89 注意 此 PDF 的内容可能过期，请检查本书网站以及 GitHub 仓库以获取最新内容。 版权声明 本书系欧长坤著，采用“知识共享署名-非商业性使用-禁止演绎 4.0 国际许可协议 (CC BY-NC-ND)”进 行许可。https://creativecommons.org/licenses/by-nc-nd/4.0/ 1 目录 目录 的正式定稿经历了长达十年多之 久的积累。C++14/17 则是作为对 C++11 的重要补充和优化，C++20 则将这门语言领进了现代化的大 门，所有这些新标准中扩充的特性，给 C++ 这门语言注入了新的活力。那些还在坚持使用传统 C++ （本书把 C++98 及其之前的 C++ 特性均称之为传统 C++）而未接触过现代 C++ 的 C++ 程序员在 见到诸如 Lambda 表达式这类全新特性时，甚至会流露出『学的不是同一门语言』的惊叹之情。 ，至少在阅读传统 C++ 代码上不具备任何困难。换句话说， 那些长期使用传统 C++ 进行编码的人、渴望在短时间内迅速了解现代 C++ 特性的人非常适合 阅读本书； 2. 本书一定程度上介绍了一些现代 C++ 的黑魔法，但这些魔法毕竟有限，不适合希望进阶学习现代 C++ 的读者，本书的定位系现代 C++ 的快速上手。当然，希望进阶学习的读者可以使用本书来 回顾并检验自己对 现代 C++ 的熟悉度。 本书目的

0.1 关于本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2 0.2 如何使用本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 4 0.3 小结 . . . . . 识，在会与 不会之间反复横跳，那么本书正是为你量身定制的！ 如果你已经积累一定的刷题量，熟悉大部分题型，那么本书可助你回顾与梳理算法知识体系，仓库源代码可 以当作“刷题工具库”或“算法字典”来使用。 若你是算法“大神”，我们期待收到你的宝贵建议，或者一起参与创作。 前置条件 你需要至少具备任一语言的编程基础，能够阅读和编写简单代码。 0.1.2 内容结构 本书的主要内容如图 0‑1 本优秀的著作。 衷心感谢我的父母，正是你们一直以来的支持与鼓励，让我有机会做这件富有趣味的事。 0.2 如何使用本书 Tip 为了获得最佳的阅读体验，建议你通读本节内容。 0.2.1 行文风格约定 ‧ 标题后标注 * 的是选读章节，内容相对困难。如果你的时间有限，可以先跳过。 ‧ 专业术语会使用黑体（纸质版和 PDF 版）或添加下划线（网页版），例如数组（array）。建议记住它们， 以便阅读文献。

0.1 关于本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2 0.2 如何使用本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 4 0.3 小结 . . . . . 识，在会与 不会之间反复横跳，那么本书正是为你量身定制的！ 如果你已经积累一定的刷题量，熟悉大部分题型，那么本书可助你回顾与梳理算法知识体系，仓库源代码可 以当作“刷题工具库”或“算法字典”来使用。 若你是算法“大神”，我们期待收到你的宝贵建议，或者一起参与创作。 � 前置条件 你需要至少具备任一语言的编程基础，能够阅读和编写简单代码。 0.1.2 内容结构 本书的主要内容如图 0‑1 机会做这件富有趣味的事。 0.2 如何使用本书 � 为了获得最佳的阅读体验，建议你通读本节内容。 0.2.1 行文风格约定 ‧ 标题后标注 * 的是选读章节，内容相对困难。如果你的时间有限，可以先跳过。 ‧ 重要专有名词及其英文翻译会用「」 括号标注，例如「数组 array」 。建议记住它们，以便阅读文献。 ‧ 专有名词和有特指含义的词句会使用“引号” 标注，以避免歧义。 ‧ 重要名词、重点内容和总结性语句会

0.1 关于本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2 0.2 如何使用本书 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 4 0.3 小结 . . . . . 认识，在会与 不会之间反复横跳，那么本书正是为你量身定制的！ 如果你已经积累一定的刷题量，熟悉大部分题型，那么本书可助你回顾与梳理算法知识体系，仓库源代码可 以当作“刷题工具库”或“算法字典”来使用。 若你是算法“大神”，我们期待收到你的宝贵建议，或者一起参与创作。 前置条件 你需要至少具备任一语言的编程基础，能够阅读和编写简单代码。 0.1.2 内容结构 本书的主要内容如图 0‑1 本优秀的著作。 衷心感谢我的父母，正是你们一直以来的支持与鼓励，让我有机会做这件富有趣味的事。 0.2 如何使用本书 Tip 为了获得最佳的阅读体验，建议你通读本节内容。 0.2.1 行文风格约定 ‧ 标题后标注 * 的是选读章节，内容相对困难。如果你的时间有限，可以先跳过。 ‧ 专业术语会使用黑体（纸质版和 PDF 版）或添加下划线（网页版），例如数组（array）。建议记住它们， 以便阅读文献。

0.1 關於本書 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 2 0.2 如何使用本書 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 4 0.3 小結 . . . . . 在會與不會之間反覆橫跳，那麼本書正是為你量身定製的！ 如果你已經積累一定的刷題量，熟悉大部分題型，那麼本書可助你回顧與梳理演算法知識體系，倉庫源程式 碼可以當作“刷題工具庫”或“演算法字典”來使用。 若你是演算法“大神”，我們期待收到你的寶貴建議，或者一起參與創作。 前置條件 你需要至少具備任一語言的程式設計基礎，能夠閱讀和編寫簡單程式碼。 0.1.2 內容結構 本書的主要內容如圖 本優秀的著作。 衷心感謝我的父母，正是你們一直以來的支持與鼓勵，讓我有機會做這件富有趣味的事。 0.2 如何使用本書 Tip 為了獲得最佳的閱讀體驗，建議你通讀本節內容。 0.2.1 行文風格約定 ‧ 標題後標註 * 的是選讀章節，內容相對困難。如果你的時間有限，可以先跳過。 ‧ 專業術語會使用黑體（紙質版和 PDF 版）或新增下劃線（網頁版），例如陣列（array）。建議記住它們， 以便閱讀文獻。