| 作者 |
| [德]贝蒂尔·施密特（Bertil Schmidt） [西]豪尔赫·冈萨雷斯-多明格斯（Jorge González-Domínguez） [德]克里斯蒂安·洪特（Christian Hundt） [德]莫里茨·施拉布（Moritz Schlarb） 等 |
| 丛书名 |
| 高性能计算技术丛书 |
| 出版社 |
| 机械工业出版社 |
| ISBN |
| 9782110261728 |
| 简要 |
| 简介 |
| 内容简介书籍计算机书籍 ---------------------------8071136 - 并行程序设计：概念与实践--------------------------- 本书由德国约翰内斯·古腾堡大学并行和分布式架构团队撰写，对并行编程的高级主题进行了深入讨论，除了涵盖基础概念外，还讲授了共享内存和分布式内存体系结构的实用编程技巧。作者提供了一套进行自动代码评估的开源系统，可方便地访问并行计算资源，这也使得本书特别适合于课堂教学。 ---------------------------8067632 - 高性能计算：现代系统与应用实践--------------------------- 本书内容全面且易学易懂，不仅涵盖高性能计算的基础概念和知识点，同时重视核心技能的培养。通过阅读本书：资深研究者将学习如何将超级计算机作为寻求新知识的重要工具；新手工程师将体验超级计算机如何运用高性能计算系统和方法来设计和模拟创新产品；学生将会了解未来高性能计算研究和开发的可能方向，为职业生涯做好准备；商用集群的维护和管理者将在掌握高性能计算系统的功能的基础上，学习如何更好地使用它们。 ---------------------------8061305 - 基于CUDA的GPU并行程序开发指南--------------------------- 本书旨在帮助读者了解与基于CUDA的并行编程技术有关的基本概念，并掌握实用c语言进行GPU高性能编程的相关技巧。本书第一部分通过CPU多线程编程解释了并行计算，使得没有太多并行计算基础的读者也能毫无阻碍地进入CUDA天地；第二部分重点介绍了基于CUDA的GPU大规模并行程序的开发与实现，并通过大量的性能分析帮助读者理解如何开发一个好的GPU并行程序以及GPU架构对程序性能的影响；本书的第三部分介绍了一些常用的CUDA库。 本书内容详尽、实例丰富，可作为高等院校相关专业高年级本科生和研究生课程的教材，也可作为有关专业研究生或计算机技术人员的参考书。 . ---------------------------5850457 - CUDA C编程权威指南--------------------------- 本书主要介绍了如何使用GPU和利用CUDAC语言对其进行编程的。首先从基本的CUDA概念及结构讲起，一步一步地引导读者进入CUDA的内部世界，由浅入深地介绍了其编程要求及其内部架构，使读者对其有了整体印象后，逐步深入了解其内部机能，后介绍了GPU的一些专用函数和注意事项。 ---------------------------5206918 - 多核与GPU编程：工具、方法及实践--------------------------- 本书从并行软件的实现、调试、优化和剖析四个方面，详细讨论了当前主要的并行计算关键技术，主要内容包括：多核和并行程序设计、共享内存编程中的线程与OpenMP、分布式内存编程、GPU编程、Thrust模板库、负载均衡等。本书结合具体的代码和案例分析，揭示了如何使用库或者指令创建多核应用，如何使用MPI开发分布式应用程序，如何使用CUDA开发高性能GPU程序，如何实现负载均衡，以及如何针对目标多核平台进行程序剖析和调试等。本书可供从事高性能计算技术研究的专业人员参考，也可作为高校相关专业的教学用书。 ---------------------------3769261 - CUDA并行程序设计：GPU编程指南--------------------------- 《CUDA并行程序设计：GPU编程指南》是CUDA并行程序设计领域最全面、最详实和最具权威性的著作之一，由CUDA开发者社区技术总监亲自撰写，英伟达中国首批CUDA官方认证工程师翻译，详实地讲解了CUDA并行程序设计的技术知识点（平台、架构、硬件知识、开发工具和热点技术）和编程方法，包含大量实用代码示例，实践性非常强。 全书共分为12章。第1章从宏观上介绍流处理器演变历史。第2章详解GPU并行机制，深入理解串行与并行程序，以辩证地求解问题。第3章讲解CUDA设备及相关的硬件和体系结构，以实现最优CUDA程序性能。第4章介绍CUDA开发环境搭建和可用调试环境。第5章介绍与CUDA编程紧密相关的核心概念——网格、线程块与线程，并通过示例说明线程模型与性能的关系。第6章借助实例详细讲解了不同类型内存的工作机制，并指出实践中容易出现的误区。第7章细述多任务的CPU和GPU协同，并介绍多个CPU/GPU编程秘技。第8章介绍如何在应用程序中编写和使用多GPU。第9章详述CUDA编程性能限制因素、分析CUDA代码的工具和技术。第10章介绍编程实践中的库与软件开发工具包。第11章讲解如何设计基于GPU的系统。第12章总结CUDA应用中易犯错误以及应对建议。 |
| 目录 |
| [套装书具体书目] 3769261 - CUDA并行程序设计：GPU编程指南 - 9787111448617 - 机械工业出版社 - 定价 99 5206918 - 多核与GPU编程：工具、方法及实践 - 9787111557685 - 机械工业出版社 - 定价 129 5850457 - CUDA C编程权威指南 - 9787111565475 - 机械工业出版社 - 定价 99 8061305 - 基于CUDA的GPU并行程序开发指南 - 9787111630616 - 机械工业出版社 - 定价 179 8067632 - 高性能计算：现代系统与应用实践 - 9787111645795 - 机械工业出版社 - 定价 149 8071136 - 并行程序设计：概念与实践 - 9787111656661 - 机械工业出版社 - 定价 119 ---------------------------8071136 - 并行程序设计：概念与实践--------------------------- 译者序 前言 致谢 第1章 绪论 1 1.1 一个有趣的例子及其分析 2 1.2 并行计算基础 10 1.2.1 分布式内存系统 10 1.2.2 共享内存系统 11 1.2.3 并行程序设计需考虑的因素 13 1.3 HPC动态和排名 15 1.4 附加练习 17 第2章 理论背景 19 2.1 PRAM 20 2.1.1 PRAM变体 21 2.1.2 PRAM上的并行前缀计算 22 2.1.3 PRAM上稀疏数组的压缩算法 24 2.2 网络拓扑 25 2.3 Amdahl定律和Gustafson定律 29 2.4 Foster的并行算法设计方法学 34 2.5 附加练习 37 参考文献 40 第3章 现代体系结构 41 3.1 存储层次 42 3.1.1 冯·诺依曼瓶颈 42 3.1.2 高速缓冲存储器 43 3.1.3 缓存算法 44 3.1.4 优化缓存访问 45 3.1.5 高速缓存一致性 48 3.1.6 虚假共享 50 3.1.7 并发多线程技术和预取技术 50 3.1.8 展望 51 3.2 并行性的层次 51 3.2.1 Flynn分类法 51 3.2.2 SIMD概念 53 3.2.3 通用微处理器上的向量化 54 3.2.4 结构体数组和数组结构体 57 3.2.5 展望 63 3.3 附加练习 63 参考文献 67 第4章 C++多线程编程 68 4.1 多线程编程简介 69 4.1.1 多线程编程和多进程编程的区别 69 4.1.2 派生和并入线程 69 4.1.3 我们的第一个多线程程序 71 4.2 处理返回值 73 4.2.1 传统方法 74 4.2.2 使用promise和future的现代方法 75 4.2.3 异步方式 80 4.3 基于静态分发的调度机制 82 4.3.1 串行程序 83 4.3.2 线程的区块分发 87 4.3.3 线程的循环分发 90 4.3.4 虚假共享 91 4.3.5 线程的块循环分发 93 4.4 处理负载不平衡 95 4.4.1 静态调度 99 4.4.2 动态块循环分发 101 4.5 用条件变量通知线程 104 4.5.1 为一个睡觉的学生建模 105 4.5.2 使用条件变量 107 4.5.3 使用future和promise单发同步 108 4.6 隐式可数集合上的并行化 110 4.6.1 隐式可数集合 111 4.6.2 线程池用例 112 4.6.3 一个简单线程池的实现 114 4.7 附加练习 119 参考文献 121 第5章 高级C++11多线程编程 122 5.1 无锁编程 122 5.1.1 原子计数 123 5.1.2 非基本原子数据类型 124 5.1.3 利用比较交换以原子方式并行化最大值归约 126 5.1.4 任意原子操作 129 5.1.5 ABA问题 132 5.2 工作共享线程池 133 5.2.1 工作共享线程池的用例 133 5.2.2 工作共享的实现 135 5.3 并行图搜索 137 5.3.1 二元背包问题 138 5.3.2 串行实现 139 5.3.3 并行实现 144 5.4 展望 146 5.5 附加练习 148 参考文献 149 第6章 OpenMP 150 6.1 OpenMP简介 151 6.1.1 OpenMP简史 151 6.1.2 基础 151 6.2 parallel for制导语句 153 6.2.1 向量加法 154 6.2.2 变量共享和私有化 157 6.2.3 矩阵向量乘法 160 6.3 基本的并行归约 162 6.3.1 最近邻分类 162 6.3.2 手写数字数据集MNIST 163 6.3.3 完全配对距离计算的理论视角 164 6.3.4 完全配对计算的实现 165 6.3.5 并行标签预测 168 6.3.6 性能评测 169 6.4 不平衡循环调度 171 6.4.1 对称性引起的负载失衡 172 6.4.2 内积计算实现 173 6.4.3 性能评测 174 6.5 高级归约 175 6.5.1 MNIST数据集上的SOFTMAX回归分类器 175 6.5.2 定制归约操作符 183 6.5.3 OpenMP高级归约 187 6.6 任务并行 189 6.6.1 树遍历 190 6.6.2 循环中生成任务 193 6.7 SIMD向量化 193 6.7.1 数据依赖 195 6.7.2 向量化感知函数 196 6.8 展望 196 6.9 附加练习 197 参考文献 202 第7章 统一计算设备架构 203 7.1 CUDA简介 204 7.2 支持CUDA的GPU硬件架构 206 7.2.1 主机与设备之间的互连 206 7.2.2 显存和峰值宽度 207 7.2.3 计算资源的组织 207 7.3 内存访问模式 211 7.3.1 均值名人脸的计算 212 7.3.2 计算中心化的数据矩阵 218 7.3.3 计算协方差矩阵 221 7.3.4 计算特征脸 229 7.4 内存层次结构 232 7.4.1 问题简介 233 7.4.2 串行DTW的线性内存算法 237 7.4.3 线性内存DTW的一个初始CUDA移植 243 7.4.4 共享内存中的波前松弛 248 7.4.5 并发调度和bank冲突 253 7.4.6 纹理内存和常量内存 254 7.5 优化准则 257 7.6 附加练习 258 参考文献 259 第8章 高级CUDA编程 261 8.1 warp内联函数和原子操作 261 8.1.1 分段并行归约 262 8.1.2 全局并行归约 265 8.1.3 任意原子操作 267 8.1.4 展望 269 8.2 利用多块GPU和流 269 8.2.1 牛顿迭代 269 8.2.2 利用多块GPU 272 8.2.3 通信和计算交叉 274 8.2.4 多块GPU上的流式计算 278 8.3 展望 280 8.3.1 统一内存 280 8.3.2 动态并行性 281 8.3.3 协作组 281 8.3.4 张量核心 281 8.3.5 GPU集群上的分布式计算 282 8.4 附加练习 282 参考文献 284 第9章 MPI 286 9.1 MPI简介 286 9.2 基本概念 288 9.3 点到点通信 289 9.4 非阻塞通信 292 9.5 集合通信 295 9.6 计算通信重叠 300 9.7 派生数据类型 309 9.8 复杂通信域 315 9.9 展望 322 9.10 附加练习 322 参考文献 327 第10章 统一并行C++ 329 10.1 PGAS和UPC++简介 329 10.2 基本概念 331 10.3 内存亲和性和私有化 332 10.4 全局指针和集合函数 337 10.5 锁 343 10.6 远程函数调用 348 10.7 附加练习 355 参考文献 357 ---------------------------8067632 - 高性能计算：现代系统与应用实践--------------------------- 译者序 序言 前言 致谢 第1章 绪论 1 1.1 高性能计算学科 2 1.1.1 定义 3 1.1.2 应用程序 3 1.1.3 性能和指标 3 1.1.4 高性能计算系统 4 1.1.5 超算问题 6 1.1.6 应用编程 7 1.2 超算对科学、社会和安全的影响 8 1.2.1 促进欺诈检测和市场数据分析 8 1.2.2 发现、管理和分配石油和天然气 8 1.2.3 加快制造业的创新 9 1.2.4 个性化医药和药物发现 9 1.2.5 预测自然灾害和了解气候变化 10 1.3 超级计算机剖析 11 1.4 计算机性能 13 1.4.1 性能 13 1.4.2 峰值性能 13 1.4.3 持续性能 14 1.4.4 可扩展性 15 1.4.5 性能退化 15 1.4.6 性能提升 17 1.5 超级计算机简史 17 1.5.1 第一个阶段—利用机械技术的自动计算机 18 1.5.2 第二个阶段—真空管时代的冯·诺依曼架构 19 1.5.3 第三个阶段—指令级并行 23 1.5.4 第四个阶段—向量处理和积分 24 1.5.5 第五个阶段—单指令多数据阵列 27 1.5.6 第六个阶段—顺序处理器的通信和超大规模集成电路 27 1.5.7 第七个阶段—多核和千万亿次 30 1.5.8 新数字时代和超越摩尔定律 30 1.6 作为学生的指南和工具 31 1.7 本章小结及成果 32 1.8 练习 33 参考文献 34 第2章 HPC架构：系统和技术 35 2.1 引言 35 2.2 HPC架构的关键特性 36 2.2.1 速度 36 2.2.2 并行性 36 2.2.3 效率 37 2.2.4 功率 37 2.2.5 可靠性 38 2.2.6 可编程性 38 2.3 并行架构家族—弗林分类法 39 2.4 支持技术 41 2.4.1 技术阶段 41 2.4.2 技术的角色 44 2.4.3 数字逻辑 45 2.4.4 存储技术 47 2.5 冯·诺依曼顺序处理器 50 2.6 向量和流水线 52 2.6.1 流水线并行 53 2.6.2 向量处理 55 2.7 单指令多数据阵列 56 2.7.1 单指令多数据架构 56 2.7.2 阿姆达定律 57 2.8 多处理器 59 2.8.1 共享内存多处理器 60 2.8.2 大规模并行处理器 62 2.8.3 商品集群系统 62 2.9 异构计算机结构 63 2.10 本章小结及成果 64 2.11 练习 65 参考文献 67 第3章 商品集群 68 3.1 引言 68 3.1.1 商品集群的定义 68 3.1.2 集群的动机和理由 69 3.1.3 集群元素 70 3.1.4 对500强名单的影响 70 3.1.5 简史 71 3.1.6 章节指南 73 3.2 Beowulf集群项目 74 3.3 硬件架构 76 3.3.1 节点 76 3.3.2 系统区域网络 77 3.3.3 辅助存储 78 3.3.4 商业系统摘要 78 3.4 编程接口 78 3.4.1 高性能计算程序设计语言 78 3.4.2 并行编程模式 80 3.5 软件环境 80 3.5.1 操作系统 80 3.5.2 资源管理 82 3.5.3 调试器 83 3.5.4 性能分析 84 3.5.5 可视化 85 3.6 基本使用方法 86 3.6.1 登录 86 3.6.2 用户空间和系统 87 3.6.3 包的配置和构建 92 3.6.4 编译器和编译 93 3.6.5 运行应用程序 94 3.7 本章小结及成果 94 3.8 练习 95 参考文献 95 第4章 基准测试程序 97 4.1 引言 97 4.2 HPC基准测试程序的关键属性 100 4.3 标准的HPC社区的基准测试程序 101 4.4 高度并行计算的Linpack 102 4.5 HPC挑战基准测试套件 104 4.6 高性能共轭梯度 106 4.7 NAS并行基准测试程序 110 4.8 Graph500 111 4.9 小型应用作为基准测试程序 115 4.10 本章小结及成果 116 4.11 练习 117 参考文献 118 第5章 资源管理的基础 119 5.1 资源管理 119 5.2 SLURM的基础 123 5.2.1 架构概述 123 5.2.2 工作负载的组织 124 5.2.3 SLURM调度 125 5.2.4 SLURM命令概要 127 5.2.5 SLURM作业脚本 141 5.2.6 SLURM速查表 146 5.3 便携式批量系统基础 147 5.3.1 PBS概述 147 5.3.2 PBS架构 148 5.3.3 PBS命令概要 148 5.3.4 PBS作业脚本 158 5.3.5 PBS速查表 160 5.4 本章小结及成果 161 5.5 练习 162 参考文献 163 第6章 对称多处理器架构 164 6.1 引言 164 6.2 架构概览 165 6.3 阿姆达定律 168 6.4 处理器核心的架构 171 6.4.1 执行流水线 172 6.4.2 指令级并行 173 6.4.3 分支预测 173 6.4.4 直通 174 6.4.5 保留站 174 6.4.6 多线程 174 6.5 存储层次 175 6.5.1 数据重用和局部性 175 6.5.2 存储层次结构 176 6.5.3 存储系统的性能 178 6.6 PCI总线 180 6.7 外部I/O接口 184 6.7.1 网络接口控制器 184 6.7.2 串行高级技术附件 185 6.7.3 JTAG 187 6.7.4 通用串行总线 189 6.8 本章小结及成果 190 6.9 练习 191 参考文献 192 第7章 OpenMP的基础 194 7.1 引言 194 7.2 OpenMP编程模型概览 195 7.2.1 线程并行 195 7.2.2 线程变量 197 7.2.3 运行时库与环境变量 197 7.3 并行线程和循环 199 7.3.1 并行线程 199 7.3.2 私有 200 7.3.3 并行“for”语句 200 7.3.4 块 206 7.4 同步 208 7.4.1 临界同步指令 208 7.4.2 master指令 209 7.4.3 barrier指令 210 7.4.4 single指令 210 7.5 归约 210 7.6 本章小结及成果 212 7.7 练习 213 参考文献 214 第8章 MPI的基础 215 8.1 引言 215 8.2 消息传递接口标准 216 8.3 消息传递接口的基本命令 217 8.3.1 mpi.h 217 8.3.2 MPI_Init 217 8.3.3 MPI_Finalize 218 8.3.4 消息传递接口例子—Hello World 218 8.4 通信器 219 8.4.1 size 220 8.4.2 rank 220 8.4.3 例子 221 8.5 点对点消息 222 8.5.1 MPI发送 222 8.5.2 消息传递接口的数据类型 222 8.5.3 MPI接收 223 8.5.4 例子 223 8.6 同步聚合 225 8.6.1 聚合调用概览 225 8.6.2 栅栏同步 225 8.6.3 例子 226 8.7 通信聚合 227 8.7.1 聚合数据移动 228 8.7.2 广播 230 8.7.3 分散 231 8.7.4 收集 233 8.7.5 全局收集 234 8.7.6 归约操作 235 8.7.7 全局到全局 238 8.8 非阻塞式点对点通信 240 8.9 用户自定义数据类型 242 8.10 本章小结及成果 243 8.11 练习 244 参考文献 244 第9章 并行算法 245 9.1 引言 245 9.2 分支-聚合 246 9.3 分治 247 9.4 管理者-工作者 249 9.5 易并行 250 9.6 光环交换 251 9.6.1 使用有限差分的平流方程 251 9.6.2 稀疏矩阵向量乘法 254 9.7 置换：佳能算法 256 9.8 任务数据流：广度优先搜索 261 9.9 本章小结及成果 263 9.10 练习 263 参考文献 264 第10章 库 265 10.1 引言 265 10.2 线性代数 267 10.2.1 基本线性代数子程序库 267 10.2.2 线性代数包 273 10.2.3 可扩展的线性代数包 275 10.2.4 GNU科学库 276 10.2.5 Supernodal LU 276 10.2.6 用于科学计算的便携式可扩展工具包 277 10.2.7 用于特征值问题计算的可扩展库 277 10.2.8 千万亿运算级的应用的特征值求解器 278 10.2.9 Hypre：可扩展的线性求解器和多重网格方法 278 10.2.10 用于线性代数的领域特定语言 278 10.3 偏微分方程 280 10.4 图算法 281 10.5 并行输入/输出 281 10.6 网格分解 288 10.7 可视化 289 10.8 并行化 289 10.9 信号处理 289 10.10 性能监控 290 10.11 本章小结及成果 290 10.12 练习 291 参考文献 292 第11章 操作系统 294 11.1 引言 294 11.2 操作系统结构和服务 296 11.2.1 系统组件 296 11.2.2 进程管理 296 11.2.3 内存管理 296 11.2.4 文件管理 296 11.2.5 I/O系统管理 297 11.2.6 辅助存储管理 297 11.3 进程管理 297 11.3.1 进程状态 298 11.3.2 进程控制块 299 11.3.3 进程管理活动 300 11.3.4 调度 301 11.4 线程 302 11.5 内存管理 304 11.5.1 虚拟内存 304 11.5.2 虚拟页地址 304 11.5.3 虚拟地址转换 305 11.6 本章小结及成果 306 11.7 练习 307 第12章 可视化 308 12.1 引言 308 12.2 基本的可视化概念 308 12.3 Gnuplot 310 12.4 Matplotlib 312 12.5 可视化工具包 316 12.6 ParaView 323 12.7 VisIt 323 12.8 本章小结及成果 324 12.9 练习 325 参考文献 325 第13章 性能监控 326 13.1 引言 326 13.2 时间测量 327 13.3 性能分析 333 13.3.1 应用分析的重要性 333 13.3.2 gperftools基础 334 13.4 监控硬件事件 339 13.4.1 perf 339 13.4.2 性能应用程序编程接口 344 13.5 集成的性能监控套件 347 13.6 分布式环境下的分析 350 13.7 本章小结及成果 356 13.8 练习 357 参考文献 358 第14章 调试 360 14.1 引言 360 14.2 工具 362 14.2.1 GNU调试器 362 14.2.2 Valgrind 370 14.2.3 商业版并行调试器 370 14.3 调试OpenMP示例：访问不受保护的共享变量 372 14.4 调试MPI示例：死锁 373 14.5 用于调试的编译器标志 378 14.6 帮助调试的系统监控器 379 14.7 本章小结及成果 381 14.8 练习 381 参考文献 385 第15章 加速器架构 386 15.1 引言 386 15.2 简史 388 15.2.1 协处理器 390 15.2.2 处理器I/O空间中的加速器 394 15.2.3 具备行业标准接口的加速器 395 15.3 图形处理单元简介 397 15.4 图形处理单元功能的进化 399 15.5 现代GPU体系结构 403 15.5.1 计算架构 403 15.5.2 内存实现 406 15.5.3 互连 407 15.5.4 可编程环境 408 15.6 异构系统架构 408 15.7 本章小结及成果 410 15.8 练习 410 参考文献 411 第16章 OpenACC的基础 412 16.1 引言 412 16.1.1 CUDA 413 16.1.2 OpenCL 414 16.1.3 C++ AMP 414 16.1.4 OpenACC 415 16.2 OpenACC编程概念 415 16.3 OpenACC库调用 417 16.4 OpenACC环境变量 419 16.5 OpenACC指令 420 16.5.1 并行构造 420 16.5.2 内核构造 422 16.5.3 数据管理 423 16.5.4 循环调度 427 16.5.5 变量的作用范围 429 16.5.6 原子性 430 16.6 本章小结及成果 431 16.7 练习 432 参考文献 433 第17章 大容量存储器 434 17.1 引言 434 17.2 存储器简史 437 17.3 存储设备技术 438 17.3.1 硬盘驱动器 438 17.3.2 固态硬盘存储器 444 17.3.3 磁带 449 17.3.4 光存储 451 17.4 集中存储 456 17.4.1 独立磁盘冗余阵列 456 17.4.2 存储区域网络 462 17.4.3 网络附加存储 463 17.4.4 三级存储器 464 17.5 本章小结及成果 465 17.6 练习 466 参考文献 466 第18章 文件系统 468 18.1 文件系统的角色和功能 468 18.2 POSIX文件接口的基础 472 18.2.1 文件访问的系统调用 472 18.2.2 缓冲文件I/O 476 18.3 网络文件系统 480 18.4 通用并行文件系统 483 18.5 Lustre文件系统 487 18.6 本章小结及成果 492 18.7 练习 493 参考文献 494 第19章 MapReduce 495 19.1 引言 495 19.2 映射和归约 495 19.2.1 单词计数 496 19.2.2 共享的邻居 497 19.2.3 K均值聚类 498 19.3 分布式计算 499 19.4 Hadoop 500 19.5 本章小结及成果 503 19.6 练习 504 参考文献 504 第20章 检查点技术 505 20.1 引言 505 20.2 系统级检查点 505 20.3 应用级检查点 511 20.4 本章小结及成果 515 20.5 练习 516 参考文献 516 第21章 下一步和未来发展 517 21.1 引言 517 21.2 扩展的并行编程模型 518 21.2.1 消息传递接口的进展 518 21.2.2 OpenMP的进展 518 21.2.3 MPI+X 519 21.3 扩展的高性能计算架构 519 21.3.1 世界上最快的机器 519 21.3.2 轻量级架构 519 21.3.3 现场可编程门阵列 520 21.4 E级计算 521 21.4.1 E级计算的挑战 522 21.4.2 从数学角度看E级的大小 523 21.4.3 加速方法 523 21.4.4 轻量级核心 523 21.5 异步多任务 524 21.5.1 多线程 524 21.5.2 消息驱动的计算 524 21.5.3 全局地址空间 525 21.5.4 行动者同步 526 21.5.5 运行时系统软件 526 21.6 新数字时代 526 21.6.1 数据流 527 21.6.2 元胞自动机 529 21.6.3 仿神经计算 530 21.6.4 量子计算 530 21.7 练习 531 参考文献 531 附录A C语言的基础 532 附录B Linux的基础 547 ---------------------------8061305 - 基于CUDA的GPU并行程序开发指南--------------------------- 译者序 前言 关于作者 第一部分 理解CPU的并行性 第1章 CPU并行编程概述 2 1.1 并行编程的演化 2 1.2 核心越多，并行性越高 3 1.3 核心与线程 4 1.3.1 并行化更多的是线程还是核心 5 1.3.2 核心资源共享的影响 6 1.3.3 内存资源共享的影响 6 1.4 第一个串行程序 7 1.4.1 理解数据传输速度 8 1.4.2 imflip.c中的main( )函数 9 1.4.3 垂直翻转行：FlipImageV( ) 10 1.4.4 水平翻转列：FlipImageH( ) 11 1.5 程序的编辑、编译、运行 12 1.5.1 选择编辑器和编译器 12 1.5.2 在Windows 7、8、10平台上开发 12 1.5.3 在Mac平台上开发 14 1.5.4 在Unix平台上开发 14 1.6 Unix速成 15 1.6.1 与相关的Unix命令 15 1.6.2 与文件相关的Unix命令 16 1.7 调试程序 19 1.7.1 gdb 19 1.7.2 古典调试方法 20 1.7.3 valgrind 22 1.8 第一个串行程序的性能 22 1.8.1 可以估计执行时间吗 23 1.8.2 代码执行时OS在做什么 23 1.8.3 如何并行化 24 1.8.4 关于资源的思考 25 第2章 开发第一个CPU并行程序 26 2.1 第一个并行程序 26 2.1.1 imflipP.c中的main( )函数 27 2.1.2 运行时间 28 2.1.3 imflipP.c中main( )函数代码的划分 28 2.1.4 线程初始化 30 2.1.5 创建线程 31 2.1.6 线程启动/执行 32 2.1.7 线程终止（合并） 33 2.1.8 线程任务和数据划分 34 2.2 位图文件 35 2.2.1 BMP是一种无损/不压缩的文件格式 35 2.2.2 BMP图像文件格式 36 2.2.3 头文件ImageStuff.h 37 2.2.4 ImageStuff.c中的图像操作函数 38 2.3 执行线程任务 40 2.3.1 启动线程 41 2.3.2 多线程垂直翻转函数MTFlipV( ) 43 2.3.3 FlipImageV( )和MTFlipV( )的比较 46 2.3.4 多线程水平翻转函数MTFlipH() 47 2.4 多线程代码的测试/计时 49 第3章 改进第一个CPU并行程序 51 3.1 程序员对性能的影响 51 3.2 CPU对性能的影响 52 3.2.1 按序核心与乱序核心 53 3.2.2 瘦线程与胖线程 55 3.3 imflipP的性能 55 3.4 操作系统对性能的影响 56 3.4.1 创建线程 57 3.4.2 线程启动和执行 57 3.4.3 线程状态 58 3.4.4 将软件线程映射到硬件线程 59 3.4.5 程序性能与启动的线程 60 3.5 改进imflipP 61 3.5.1 分析MTFlipH( )中的内存访问模式 62 3.5.2 MTFlipH( )的多线程内存访问 63 3.5.3 DRAM访问的规则 64 3.6 imflipPM：遵循DRAM的规则 65 3.6.1 imflipP的混乱内存访问模式 65 3.6.2 改进imflipP的内存访问模式 65 3.6.3 MTFlipHM( )：内存友好的MTFlipH( ) 66 3.6.4 MTFlipVM( )：内存友好的MTFlipV( ) 69 3.7 imflipPM.C的性能 69 3.7.1 imflipP.c和imflipPM.c的性能比较 70 3.7.2 速度提升：MTFlipV( )与MTFlipVM( ) 71 3.7.3 速度提升：MTFlipH( )与MTFlipHM( ) 71 3.7.4 理解加速：MTFlipH( )与MTFlipHM( ) 71 3.8 进程内存映像 72 3.9 英特尔MIC架构：Xeon Phi 74 3.10 GPU是怎样的 75 3.11 本章小结 76 第4章 理解核心和内存 77 4.1 曾经的英特尔 77 4.2 CPU和内存制造商 78 4.3 动态存储器与静态存储器 79 4.3.1 静态随机存取存储器（SRAM） 79 4.3.2 动态随机存取存储器（DRAM） 79 4.3.3 DRAM接口标准 79 4.3.4 DRAM对程序性能的影响 80 4.3.5 SRAM对程序性能的影响 81 4.4 图像旋转程序：imrotate.c 81 4.4.1 imrotate.c的说明 82 4.4.2 imrotate.c：参数限制和简化 82 4.4.3 imrotate.c：实现原理 83 4.5 imrotate的性能 87 4.5.1 线程效率的定性分析 87 4.5.2 定量分析：定义线程效率 87 4.6 计算机的体系结构 89 4.6.1 核心、L1$和L2$ 89 4.6.2 核心内部资源 90 4.6.3 共享L3高速缓存（L3 $） 91 4.6.4 内存控制器 92 4.6.5 主存 92 4.6.6 队列、非核心和I/O 93 4.7 imrotateMC：让imrotate更高效 94 4.7.1 Rotate2( )：平方根和浮点除法有多差 96 4.7.2 Rotate3( )和Rotate4( )：sin( )和cos( )有多差 97 4.7.3 Rotate5( )：整数除法/乘法有多差 98 4.7.4 Rotate6( )：合并计算 100 4.7.5 Rotate7( )：合并更多计算 100 4.7.6 imrotateMC的总体性能 101 4.8 本章小结 103 第5章 线程管理和同步 104 5.1 边缘检测程序：imedge.c 104 5.1.1 imedge.c的说明 105 5.1.2 imedge.c：参数限制和简化 106 5.1.3 imedge.c：实现原理 106 5.2 imedge.c：实现 108 5.2.1 初始化和时间戳 109 5.2.2 不同图像表示的初始化函数 110 5.2.3 启动和终止线程 111 5.2.4 高斯滤波 112 5.2.5 Sobel 113 5.2.6 阈值过滤 114 5.3 imedge的性能 115 5.4 imedgeMC：让imedge更高效 116 5.4.1 利用预计算降低带宽 116 5.4.2 存储预计算的像素值 117 5.4.3 预计算像素值 118 5.4.4 读取图像并预计算像素值 119 5.4.5 PrGaussianFilter 120 5.4.6 PrSobel 121 5.4.7 PrThreshold 121 5.5 imedgeMC的性能 122 5.6 imedgeMCT：高效的线程同步 123 5.6.1 屏障同步 124 5.6.2 用于数据共享的MUTEX结构 125 5.7 imedgeMCT：实现 127 5.7.1 使用MUTEX：读取图像、预计算 128 5.7.2 一次预计算一行 130 5.8 imedgeMCT的性能 131 第二部分 基于CUDA的GPU编程 第6章 GPU并行性和CUDA概述 134 6.1 曾经的Nvidia 134 6.1.1 GPU的诞生 134 6.1.2 早期的GPU架构 136 6.1.3 GPGPU的诞生 137 6.1.4 Nvidia、ATI Technologies和Intel 138 6.2 统一计算设备架构 140 6.2.1 CUDA、OpenCL和其他GPU语言 140 6.2.2 设备端与主机端代码 140 6.3 理解GPU并行 141 6.3.1 GPU如何实现高性能 142 6.3.2 CPU与GPU架构的差异 143 6.4 图像翻转的CUDA版：imflipG.cu 144 6.4.1 imflipG.cu：将图像读入CPU端数组 146 6.4.2 初始化和查询GPU 147 6.4.3 GPU端的时间戳 148 6.4.4 GPU端内存分配 152 6.4.5 GPU驱动程序和Nvidia运行时引擎 153 6.4.6 CPU到GPU的数据传输 153 6.4.7 用封装函数进行错误报告 154 6.4.8 GPU核函数的执行 155 6.4.9 完成GPU核函数的执行 157 6.4.10 将GPU结果传回CPU 158 6.4.11 完成时间戳 158 6.4.12 输出结果和清理 158 6.4.13 读取和输出BMP文件 159 6.4.14 Vflip( )：垂直翻转的GPU核函数 160 6.4.15 什么是线程ID、块ID和块维度 163 6.4.16 Hflip( )：水平翻转的GPU核函数 165 6.4.17 硬件参数：threadIDx.x、blockIdx.x和blockDim.x 165 6.4.18 PixCopy( )：复制图像的GPU核函数 165 6.4.19 CUDA关键字 166 6.5 Windows中的CUDA程序开发 167 6.5.1 安装MS Visual Studio 2015和CUDA Toolkit 8.0 167 6.5.2 在Visual Studio 2015中创建项目imflipG.cu 168 6.5.3 在Visual Studio 2015中编译项目imflipG.cu 170 6.5.4 运行第一个CUDA应用程序：imflipG.exe 173 6.5.5 确保程序的正确性 174 6.6 Mac平台上的CUDA程序开发 175 6.6.1 在Mac上安装XCode 175 6.6.2 安装CUDA驱动程序和CUDA工具包 176 6.6.3 在Mac上编译和运行CUDA应用程序 177 6.7 Unix平台上的CUDA程序开发 177 6.7.1 安装Eclipse和CUDA工具包 177 6.7.2 使用ssh登录一个集群 178 6.7.3 编译和执行CUDA代码 179 第7章 CUDA主机/设备编程模型 181 7.1 设计程序的并行性 181 7.1.1 任务的并行化 182 7.1.2 什么是Vflip( )的最佳块尺寸 183 7.1.3 imflipG.cu：程序输出的解释 183 7.1.4 imflipG.cu：线程块和图像的大小对性能的影响 184 7.2 核函数的启动 185 7.2.1 网格 185 7.2.2 线程块 187 7.2.3 线程 187 7.2.4 线程束和通道 189 7.3 imflipG.cu：理解核函数的细节 189 7.3.1 在main( )中启动核函数并将参数传递给它们 189 7.3.2 线程执行步骤 190 7.3.3 Vflip( )核函数 191 7.3.4 Vflip( )和MTFlipV( )的比较 192 7.3.5 Hflip( )核函数 194 7.3.6 PixCopy( )核函数 194 7.4 PCIe速度与CPU的关系 196 7.5 PCIe总线对性能的影响 196 7.5.1 数据传输时间、速度、延迟、吞吐量和带宽 196 7.5.2 imflipG.cu的PCIe吞吐量 197 7.6 全局内存总线对性能的影响 200 7.7 计算能力对性能的影响 203 7.7.1 Fermi、Kepler、Maxwell、Pascal和Volta系列 203 7.7.2 不同系列实现的相对带宽 204 7.7.3 imflipG2.cu：计算能力2.0版本的imflipG.cu 205 7.7.4 imflipG2.cu：main( )的修改 206 7.7.5 核函数PxCC20( ) 208 7.7.6 核函数VfCC20( ) 208 7.8 imflipG2.cu的性能 210 7.9 古典的CUDA调试方法 212 7.9.1 常见的CUDA错误 212 7.9.2 return调试法 214 7.9.3 基于注释的调试 216 7.9.4 printf()调试 216 7.10 软件错误的生物学原因 217 7.10.1 大脑如何参与编写/调试代码 218 7.10.2 当我们疲倦时是否会写出错误代码 219 第8章 理解GPU的硬件架构 221 8.1 GPU硬件架构 222 8.2 GPU硬件的部件 222 8.2.1 SM：流处理器 222 8.2.2 GPU核心 223 8.2.3 千兆线程调度器 223 8.2.4 内存控制器 225 8.2.5 共享高速缓存（L2$） 225 8.2.6 主机接口 225 8.3 Nvidia GPU架构 226 8.3.1 Fermi架构 227 8.3.2 GT、GTX和计算加速器 227 8.3.3 Kepler架构 228 8.3.4 Maxwell架构 228 8.3.5 Pascal架构和NVLink 229 8.4 CUDA边缘检测：imedgeG.cu 229 8.4.1 CPU和GPU内存中存储图像的变量 229 8.4.2 为GPU变量分配内存 231 8.4.3 调用核函数并对其进行计时 233 8.4.4 计算核函数的性能 234 8.4.5 计算核函数的数据移动量 235 8.4.6 输出核函数性能 236 8.5 imedgeG：核函数 237 8.5.1 BWKernel( ) 237 8.5.2 GaussKernel( ) 239 8.5.3 SobelKernel( ) 240 8.5.4 ThresholdKernel( ) 242 8.6 imedgeG.cu的性能 243 8.6.1 imedgeG.cu：PCIe总线利用率 244 8.6.2 imedgeG.cu：运行时间 245 8.6.3 imedgeG.cu：核函数性能比较 247 8.7 GPU代码：编译时间 248 8.7.1 设计CUDA代码 248 8.7.2 编译CUDA代码 250 8.7.3 GPU汇编：PTX、CUBIN 250 8.8 GPU代码：启动 250 8.8.1 操作系统的参与和CUDA DLL文件 250 8.8.2 GPU图形驱动 251 8.8.3 CPU与GPU之间的内存传输 251 8.9 GPU代码：执行（运行时间） 252 8.9.1 获取数据 252 8.9.2 获取代码和参数 252 8.9.3 启动线程块网格 252 8.9.4 千兆线程调度器（GTS） 253 8.9.5 线程块调度 254 8.9.6 线程块的执行 255 8.9.7 透明的可扩展性 256 第9章 理解GPU核心 257 9.1 GPU的架构系列 258 9.1.1 Fermi架构 258 9.1.2 Fermi SM的结构 259 9.1.3 Kepler架构 260 9.1.4 Kepler SMX的结构 260 9.1.5 Maxwell架构 261 9.1.6 Maxwell SMM的结构 262 9.1.7 Pascal GP100架构 264 9.1.8 Pascal GP100 SM的结构 265 9.1.9 系列比较：峰值GFLOPS和峰值DGFLOPS 266 9.1.10 GPU睿频 267 9.1.11 GPU功耗 268 9.1.12 计算机电源 268 9.2 流处理器的构建模块 269 9.2.1 GPU核心 269 9.2.2 双精度单元（DPU） 270 9.2.3 特殊功能单元（SFU） 270 9.2.4 寄存器文件（RF） 270 9.2.5 读取/存储队列（LDST） 271 9.2.6 L1$和纹理高速缓存 272 9.2.7 共享内存 272 9.2.8 常量高速缓存 272 9.2.9 指令高速缓存 272 9.2.10 指令缓冲区 272 9.2.11 线程束调度器 272 9.2.12 分发单元 273 9.3 并行线程执行（PTX）的数据类型 273 9.3.1 INT8：8位整数 274 9.3.2 INT16：16位整数 274 9.3.3 24位整数 275 9.3.4 INT32：32位整数 275 9.3.5 判定寄存器（32位） 275 9.3.6 INT64：64位整数 276 9.3.7 128位整数 276 9.3.8 FP32：单精度浮点（float） 277 9.3.9 FP64：双精度浮点（double） 277 9.3.10 FP16：半精度浮点（half） 278 9.3.11 什么是FLOP 278 9.3.12 融合乘法累加（FMA）与乘加（MAD） 278 9.3.13 四倍和八倍精度浮点 279 9.3.14 Pascal GP104引擎的SM结构 279 9.4 imflipGC.cu：核心友好的imflipG 280 9.4.1 Hflip2( )：预计算核函数参数 282 9.4.2 Vflip2( )：预计算核函数参数 284 9.4.3 使用线程计算图像坐标 285 9.4.4 线程块ID与图像的行映射 285 9.4.5 Hflip3( )：使用二维启动网格 286 9.4.6 Vflip3( )：使用二维启动网格 287 9.4.7 Hflip4( )：计算2个连续的像素 288 9.4.8 Vflip4( )：计算2个连续的像素 289 9.4.9 Hflip5( )：计算4个连续的像素 290 9.4.10 Vflip5( )：计算4个连续的像素 291 9.4.11 PixCopy2( )、PixCopy3( )：一次分别复制2个和4个连续的像素 292 9.5 imedgeGC.cu：核心友好的imedgeG 293 9.5.1 BWKernel2( )：使用预计算的值和2D块 293 9.5.2 GaussKernel2( )：使用预计算的值和2D块 294 第10章 理解GPU内存 296 10.1 全局内存 297 10.2 L2高速缓存 297 10.3 纹理/L1高速缓存 298 10.4 共享内存 298 10.4.1 分拆与专用共享内存 299 10.4.2 每核心可用的内存资源 299 10.4.3 使用共享内存作为软件高速缓存 300 10.4.4 分配SM中的共享内存 300 10.5 指令高速缓存 300 10.6 常量内存 301 10.7 imflipGCM.cu：核心和内存友好的imflipG 301 10.7.1 Hflip6( )、Vflip6( )：使用共享内存作为缓冲区 301 10.7.2 Hflip7( )：共享内存中连续的交换操作 303 10.7.3 Hflip8( )：使用寄存器交换4个像素 305 10.7.4 Vflip7( )：一次复制4个字节（int） 307 10.7.5 对齐与未对齐的内存数据访问 308 10.7.6 Vflip8( )：一次复制8个字节 308 10.7.7 Vflip9( )：仅使用全局内存，一次复制8个字节 309 10.7.8 PixCopy4( )、PixCopy5( )：使用共享内存复制1个和4个字节 310 10.7.9 PixCopy6( )、PixCopy7( )：使用全局内存复制1个和2个整数 311 10.8 imedgeGCM.cu：核心和内存友好的imedgeG 312 10.8.1 BWKernel3( )：使用字节操作来提取RGB 312 10.8.2 GaussKernel3( )：使用常量内存 314 10.8.3 处理常量的方法 315 10.8.4 GaussKernel4( )：在共享内存中缓冲1个像素的邻居 316 10.8.5 GaussKernel5( )：在共享内存中缓冲4个像素的邻居 318 10.8.6 GaussKernel6( )：将5个垂直像素读入共享内存 320 10.8.7 GaussKernel7( )：去除边界像素的影响 322 10.8.8 GaussKernel8( )：计算8个垂直像素 324 10.9 CUDA占用率计算器 326 10.9.1 选择最佳的线程/块 327 10.9.2 SM级资源限制 328 10.9.3 什么是占用率 329 10.9.4 CUDA占用率计算器：资源计算 330 10.9.5 案例分析：GaussKernel7() 334 10.9.6 案例分析：GaussKernel8() 337 第11章 CUDA流 340 11.1 什么是流水线 342 11.1.1 重叠执行 342 11.1.2 暴露与合并的运行时间 343 11.2 内存分配 344 11.2.1 物理与虚拟内存 345 11.2.2 物理地址到虚拟地址的转换 345 11.2.3 固定内存 345 11.2.4 使用cudaMallocHost( )分配固定内存 346 11.3 CPU与GPU之间快速的数据传输 346 11.3.1 同步数据传输 346 11.3.2 异步数据传输 347 11.4 CUDA流的原理 347 11.4.1 CPU到GPU的传输、核函数的执行、GPU到CPU的传输 347 11.4.2 在CUDA中实现流 348 11.4.3 复制引擎 348 11.4.4 核函数执行引擎 349 11.4.5 并发的上行和下行PCIe传输 349 11.4.6 创建CUDA流 350 11.4.7 销毁CUDA流 350 11.4.8 同步CUDA流 350 11.5 imGStr.cu：流式图像处理 351 11.5.1 将图像读入固定内存 351 11.5.2 同步与单个流 353 11.5.3 多个流 354 11.5.4 多流之间的数据依赖 356 11.6 流式水平翻转核函数 360 11.7 imGStr.cu：流式边缘检测 362 11.8 性能对比：imGStr.cu 365 11.8.1 同步与异步结果 366 11.8.2 结果的随机性 366 11.8.3 队列优化 366 11.8.4 最佳流式结果 367 11.8.5 最差流式结果 369 11.9 Nvidia可视化分析器：nvvp 370 11.9.1 安装nvvp和nvprof 370 11.9.2 使用nvvp 370 11.9.3 使用nvprof 371 11.9.4 imGStr的同步和单流结果 372 11.9.5 imGStr的2流和4流结果 373 第三部分 拓展知识 第12章 CUDA库 376 12.1 cuBLAS 377 12.1.1 BLAS级别 377 12.1.2 cuBLAS数据类型 377 12.1.3 安装cuBLAS 378 12.1.4 变量声明和初始化 378 12.1.5 设备内存分配 379 12.1.6 创建上下文 379 12.1.7 将数据传输到设备端 379 12.1.8 调用cuBLAS函数 380 12.1.9 将数据传回主机 380 12.1.10 释放内存 381 12.1.11 cuBLAS程序示例：矩阵的标量操作 381 12.2 cuFFT 382 12.2.1 cuFFT库特征 383 12.2.2 复数到复数变换示例 383 12.2.3 实数到复数变换示例 384 12.3 Nvidia性能库（NPP） 384 12.4 Thrust库 386 第13章 OpenCL简介 388 13.1 什么是OpenCL 388 13.1.1 多平台 388 13.1.2 基于队列 389 13.2 图像翻转核函数 389 13.3 运行核函数 390 13.3.1 选择设备 390 13.3.2 运行核函数 392 13.3.3 OpenCL程序的运行时间 396 13.4 OpenCL中的边缘检测 396 第14章 其他GPU编程语言 402 14.1 使用Python进行GPU编程 402 14.1.1 imflip的PyOpenCL版本 403 14.1.2 PyOpenCL的逐元素核函数 406 14.2 OpenGL 408 14.3 OpenGL ES：用于嵌入式系统的OpenGL 409 14.4 Vulkan 409 14.5 微软的高级着色语言 409 14.5.1 着色 410 14.5.2 HLSL 410 14.6 Apple的Metal API 411 14.7 Apple的Swift编程语言 411 14.8 OpenCV 411 14.8.1 安装OpenCV和人脸识别 411 14.8.2 移动–微云–云实时人脸识别 412 14.8.3 加速即服务（AXaas） 412 第15章 深度学习中的CUDA 413 15.1 人工神经网络 413 15.1.1 神经元 413 15.1.2 激活函数 414 15.2 全连接神经网络 415 15.3 深度网络/卷积神经网络 415 15.4 训练网络 416 15.5 cuDNN深度学习库 416 15.5.1 创建一个层 417 15.5.2 创建一个网络 418 15.5.3 前向传播 418 15.5.4 反向传播 419 15.5.5 在网络中使用cuBLAS 419 15.6 Keras 419 参考文献 421 术语表 424 ---------------------------5850457 - CUDA C编程权威指南--------------------------- 译者序 推荐序 自序 作者简介 技术审校者简介 前言 致谢 第1章 基于CUDA的异构并行计算1 1.1 并行计算1 1.1.1 串行编程和并行编程2 1.1.2 并行性3 1.1.3 计算机架构4 1.2 异构计算6 1.2.1 异构架构7 1.2.2 异构计算范例9 1.2.3 CUDA：一种异构计算平台10 1.3 用GPU输出Hello World12 1.4 使用CUDA C编程难吗15 1.5 总结16 1.6 习题16 第2章 CUDA编程模型18 2.1 CUDA编程模型概述18 2.1.1 CUDA编程结构19 2.1.2 内存管理20 2.1.3 线程管理24 2.1.4 启动一个CUDA核函数29 2.1.5 编写核函数30 2.1.6 验证核函数31 2.1.7 处理错误32 2.1.8 编译和执行32 2.2 给核函数计时35 2.2.1 用CPU计时器计时35 2.2.2 用nvprof工具计时39 2.3 组织并行线程40 2.3.1 使用块和线程建立矩阵索引40 2.3.2 使用二维网格和二维块对矩阵求和44 2.3.3 使用一维网格和一维块对矩阵求和47 2.3.4 使用二维网格和一维块对矩阵求和48 2.4 设备管理50 2.4.1 使用运行时API查询GPU信息50 2.4.2 确定最优GPU53 2.4.3 使用nvidia-smi查询GPU信息53 2.4.4 在运行时设置设备54 2.5 总结54 2.6 习题55 第3章 CUDA执行模型56 3.1 CUDA执行模型概述56 3.1.1 GPU架构概述57 3.1.2 Fermi架构59 3.1.3 Kepler架构61 3.1.4 配置文件驱动优化65 3.2 理解线程束执行的本质67 3.2.1 线程束和线程块67 3.2.2 线程束分化69 3.2.3 资源分配74 3.2.4 延迟隐藏76 3.2.5 占用率78 3.2.6 同步81 3.2.7 可扩展性82 3.3 并行性的表现83 3.3.1 用nvprof检测活跃的线程束84 3.3.2 用nvprof检测内存操作85 3.3.3 增大并行性86 3.4 避免分支分化88 3.4.1 并行归约问题88 3.4.2 并行归约中的分化89 3.4.3 改善并行归约的分化93 3.4.4 交错配对的归约95 3.5 展开循环97 3.5.1 展开的归约97 3.5.2 展开线程的归约99 3.5.3 完全展开的归约101 3.5.4 模板函数的归约102 3.6 动态并行104 3.6.1 嵌套执行105 3.6.2 在GPU上嵌套Hello World106 3.6.3 嵌套归约109 3.7 总结113 3.8 习题113 第4章 全局内存115 4.1 CUDA内存模型概述115 4.1.1 内存层次结构的优点116 4.1.2 CUDA内存模型117 4.2 内存管理124 4.2.1 内存分配和释放124 4.2.2 内存传输125 4.2.3 固定内存127 4.2.4 零拷贝内存128 4.2.5 统一虚拟寻址133 4.2.6 统一内存寻址134 4.3 内存访问模式135 4.3.1 对齐与合并访问135 4.3.2 全局内存读取137 4.3.3 全局内存写入145 4.3.4 结构体数组与数组结构体147 4.3.5 性能调整151 4.4 核函数可达到的带宽154 4.4.1 内存带宽154 4.4.2 矩阵转置问题155 4.5 使用统一内存的矩阵加法167 4.6 总结171 4.7 习题172 第5章 共享内存和常量内存174 5.1 CUDA共享内存概述174 5.1.1 共享内存175 5.1.2 共享内存分配176 5.1.3 共享内存存储体和访问模式176 5.1.4 配置共享内存量181 5.1.5 同步183 5.2 共享内存的数据布局185 5.2.1 方形共享内存185 5.2.2 矩形共享内存193 5.3 减少全局内存访问199 5.3.1 使用共享内存的并行归约199 5.3.2 使用展开的并行归约202 5.3.3 使用动态共享内存的并行归约204 5.3.4 有效带宽205 5.4 合并的全局内存访问205 5.4.1 基准转置内核205 5.4.2 使用共享内存的矩阵转置207 5.4.3 使用填充共享内存的矩阵转置210 5.4.4 使用展开的矩阵转置211 5.4.5 增大并行性214 5.5 常量内存215 5.5.1 使用常量内存实现一维模板215 5.5.2 与只读缓存的比较217 5.6 线程束洗牌指令219 5.6.1 线程束洗牌指令的不同形式220 5.6.2 线程束内的共享数据222 5.6.3 使用线程束洗牌指令的并行归约226 5.7 总结227 5.8 习题228 第6章 流和并发230 6.1 流和事件概述231 6.1.1 CUDA流231 6.1.2 流调度234 6.1.3 流的优先级235 6.1.4 CUDA事件235 6.1.5 流同步237 6.2 并发内核执行240 6.2.1 非空流中的并发内核240 6.2.2 Fermi GPU上的虚假依赖关系242 6.2.3 使用OpenMP的调度操作244 6.2.4 用环境变量调整流行为245 6.2.5 GPU资源的并发限制246 6.2.6 默认流的阻塞行为247 6.2.7 创建流间依赖关系248 6.3 重叠内核执行和数据传输249 6.3.1 使用深度优先调度重叠249 6.3.2 使用广度优先调度重叠252 6.4 重叠GPU和CPU执行254 6.5 流回调255 6.6 总结256 6.7 习题257 第7章 调整指令级原语258 7.1 CUDA指令概述259 7.1.1 浮点指令259 7.1.2 内部函数和标准函数261 7.1.3 原子操作指令262 7.2 程序优化指令264 7.2.1 单精度与双精度的比较264 7.2.2 标准函数与内部函数的比较266 7.2.3 了解原子指令272 7.2.4 综合范例277 7.3 总结279 7.4 习题280 第8章 GPU加速库和OpenACC281 8.1 CUDA库概述282 8.1.1 CUDA库支持的作用域283 8.1.2 通用的CUDA库工作流283 8.2 cuSPARSE库285 8.2.1 cuSPARSE数据存储格式286 8.2.2 用cuSPARSE进行格式转换289 8.2.3 cuSPARSE功能示例289 8.2.4 cuSPARSE发展中的重要主题291 8.2.5 cuSPARSE小结291 8.3 cuBLAS库292 8.3.1 管理cuBLAS数据293 8.3.2 cuBLAS功能示例294 8.3.3 cuBLAS发展中的重要主题295 8.3.4 cuBLAS小结296 8.4 cuFFT库296 8.4.1 使用cuFFT API296 8.4.2 cuFFT功能示例298 8.4.3 cuFFT小结299 8.5 cuRAND库299 8.5.1 拟随机数或伪随机数的选择299 8.5.2 cuRAND库概述300 8.5.3 cuRAND介绍303 8.5.4 cuRAND发展中的重要主题306 8.6 CUDA 6.0中函数库的介绍307 8.6.1 Drop-In库307 8.6.2 多GPU库308 8.7 CUDA函数库的性能研究310 8.7.1 cuSPARSE与MKL的比较310 8.7.2 cuBLAS与MKL BLAS的比较311 8.7.3 cuFFT与FFTW及MKL的比较311 8.7.4 CUDA库性能小结312 8.8 OpenACC的使用312 8.8.1 OpenACC计算指令的使用315 8.8.2 OpenACC数据指令的使用321 8.8.3 OpenACC运行时API325 8.8.4 OpenACC和CUDA库的结合327 8.8.5 OpenACC小结328 8.9 总结329 8.10 习题329 第9章 多GPU编程331 9.1 从一个GPU到多GPU332 9.1.1 在多GPU上执行333 9.1.2 点对点通信334 9.1.3 多GPU间的同步335 9.2 多GPU间细分计算336 9.2.1 在多设备上分配内存336 9.2.2 单主机线程分配工作337 9.2.3 编译和执行337 9.3 多GPU上的点对点通信338 9.3.1 实现点对点访问338 9.3.2 点对点的内存复制339 9.3.3 统一虚拟寻址的点对点内存访问341 9.4 多GPU上的有限差分342 9.4.1 二维波动方程的模板计算342 9.4.2 多GPU程序的典型模式343 9.4.3 多GPU上的二维模板计算344 9.4.4 重叠计算与通信347 9.4.5 编译和执行348 9.5 跨GPU集群扩展应用程序350 9.5.1 CPU到CPU的数据传输351 9.5.2 使用传统MPI在GPU和GPU间传输数据353 9.5.3 使用CUDA-aware MPI进行GPU到GPU的数据传输356 9.5.4 使用CUDA-aware MPI进行节点内GPU到GPU的数据传输357 9.5.5 调整消息块大小358 9.5.6 使用GPUDirect RDMA技术进行GPU到GPU的数据传输359 9.6 总结361 9.7 习题362 第10章 程序实现的注意事项364 10.1 CUDA C的开发过程364 10.1.1 APOD开发周期365 10.1.2 优化因素367 10.1.3 CUDA代码编译370 10.1.4 CUDA错误处理373 10.2 配置文件驱动优化374 10.2.1 使用nvprof寻找优化因素375 10.2.2 使用nvvp指导优化379 10.2.3 NVIDIA工具扩展381 10.3 CUDA调试383 10.3.1 内核调试383 10.3.2 内存调试390 10.3.3 调试小结395 10.4 将C程序移植到CUDA C的案例研究396 10.4.1 评估crypt396 10.4.2 并行crypt397 10.4.3 优化crypt398 10.4.4 部署crypt404 10.4.5 移植crypt小结407 10.5 总结407 10.6 习题407 附录 推荐阅读409 ---------------------------5206918 - 多核与GPU编程：工具、方法及实践--------------------------- 译者序 前 言 第1章 概述 1 1.1 多核计算机时代 1 1.2 并行计算机的分类 3 1.3 现代计算机概览 4 1.3.1 Cell BE处理器 5 1.3.2 NVIDIA Kepler 6 1.3.3 AMD APU 9 1.3.4 从多核到众核：Tilera TILE-Gx8072和Intel Xeon Phi 10 1.4 性能指标 12 1.5 并行程序性能的预测与测量 16 1.5.1 Amdahl定律 18 1.5.2 Gustafson-Barsis定律 20 第2章 多核和并行程序设计 23 2.1 引言 23 2.2 PCAM方法学 24 2.3 分解模式 26 2.3.1 任务并行 27 2.3.2 分而治之分解 28 2.3.3 几何分解 30 2.3.4 递归数据分解 32 2.3.5 流水线分解 35 2.3.6 基于事件的合作分解 39 2.4 程序结构模式 39 2.4.1 单程序多数据 40 2.4.2 多程序多数据 40 2.4.3 主/从 41 2.4.4 map-reduce 41 2.4.5 fork/join 42 2.4.6 循环并行 44 2.5 匹配分解模式和程序结构模式 44 第3章 共享内存编程：线程 46 3.1 引言 46 3.2 线程 48 3.2.1 线程的定义 48 3.2.2 线程的作用 49 3.2.3 线程的生成和初始化 49 3.2.4 在线程间共享数据 55 3.3 设计考虑 57 3.4 信号量 58 3.5 经典问题中的信号量 62 3.5.1 生产者–消费者 63 3.5.2 终止处理 66 3.5.3 理发师问题：引入公平性 75 3.5.4 读者–写者问题 80 3.6 monitor 84 3.6.1 设计方法1：monitor内部的关键区 87 3.6.2 设计方法2：monitor控制关键区的入口 87 3.7 经典问题中的monitor 91 3.7.1 重新考虑生产者–消费者问题 91 3.7.2 重新考虑读者–写者问题 95 3.8 动态线程管理与静态线程管理 102 3.8.1 Qt线程池 102 3.8.2 线程池的创建和管理 103 3.9 调试多线程应用 111 3.10 高层次结构：无须显式利用线程的多线程编程 115 3.10.1 并发map 116 3.10.2 map-reduce 118 3.10.3 并发过滤 120 3.10.4 filter-reduce 121 3.10.5 案例研究：多线程存储 122 3.10.6 案例研究：多线程图像匹配 131 第4章 共享内存编程：OpenMP 140 4.1 引言 140 4.2 第一个OpenMP程序 141 4.3 变量作用域 144 4.3.1 定积分OpenMP版本V.0：人工划分 146 4.3.2 定积分OpenMP版本 V.1：无竞争条件的人工划分 147 4.3.3 定积分OpenMP V.2：基于锁的隐式划分 148 4.3.4 定积分OpenMP V.3：基于归约的隐式划分 150 4.3.5 变量作用域总结 151 4.4 循环级并行 152 4.4.1 数据依赖 154 4.4.2 嵌套循环 162 4.4.3 调度 162 4.5 任务并行 166 4.5.1 sections指令 166 4.5.2 task指令 171 4.6 同步结构 177 4.7 正确性与优化问题 183 4.7.1 线程安全 183 4.7.2 假共享 187 4.8 案例研究：OpenMP中的排序算法 192 4.8.1 自下而上归并排序算法的OpenMP实现 192 4.8.2 自上而下归并排序算法的OpenMP实现 195 4.8.3 性能评估 200 第5章 分布式内存编程 203 5.1 通信进程 203 5.2 MPI 204 5.3 核心概念 205 5.4 你的第一个MPI程序 206 5.5 程序体系结构 208 5.5.1 SPMD 208 5.5.2 MPMD 209 5.6 点对点通信 210 5.7 可选的点对点通信模式 214 5.8 非阻塞通信 216 5.9 点对点通信小结 220 5.10 错误报告与处理 220 5.11 集合通信简介 222 5.11.1 分发 226 5.11.2 收集 231 5.11.3 归约 233 5.11.4 多对多收集 237 5.11.5 多对多分发 240 5.11.6 多对多归约 245 5.11.7 全局同步 245 5.12 通信对象 245 5.12.1 派生数据类型 246 5.12.2 打包/解包 253 5.13 节点管理：通信器和组 254 5.13.1 创建组 255 5.13.2 建立内部通信器 257 5.14 单边通信 259 5.14.1 RMA通信函数 261 5.14.2 RMA同步函数 262 5.15 I/O注意事项 270 5.16 MPI多进程和多线程混合编程 276 5.17 时序和性能测量 279 5.18 调试和分析MPI程序 279 5.19 Boost.MPI库 283 5.19.1 阻塞和非阻塞通信 285 5.19.2数据序列化 289 5.19.3集合通信 292 5.20 案例研究：有限扩散聚合模型 295 5.21 案例研究：暴力加密破解 300 5.21.1 版本1：“基本型”MPI 300 5.21.2 版本2：MPI与OpenMP的结合 305 5.22 案例研究：主/从式并行模型的MPI实现 308 5.22.1 简单主/从式设置 309 5.22.2 多线程主/从式设置 316 第6章 GPU编程 333 6.1 GPU编程简介 333 6.2 CUDA编程模型：线程、线程块、线程网格 335 6.3 CUDA执行模型：流多处理器和warp 340 6.4 CUDA程序编译过程 344 6.5 构建CUDA项目 347 6.6 内存层次结构 349 6.6.1 本地内存/寄存器 355 6.6.2 共享内存 356 6.6.3 常量内存 363 6.6.4 texture和surface内存 368 6.7 优化技术 369 6.7.1 线程组织设计 369 6.7.2 kernel结构 378 6.7.3 共享内存访问 382 6.7.4 全局内存访问 388 6.7.5 page-locked与zero-copy内存 392 6.7.6 统一内存 394 6.7.7 异步执行和流 397 6.8 动态并行 403 6.9 CUDA程序的调试 407 6.10 CUDA程序剖析 410 6.11 CUDA和MPI 412 6.12 案例研究 417 6.12.1 分形集合计算 417 6.12.2 块加密算法 426 第7章 Thrust模板库 452 7.1 引言 452 7.2 使用Thrust的第一步 453 7.3 Thrust数据类型 456 7.4 Thrust算法 459 7.4.1 变换算法 460 7.4.2 排序与查询 463 7.4.3 归约 468 7.4.4 scan /前缀和 471 7.4.5 数据管理与处理 472 7.5 花式迭代器 475 7.6 交换设备后端 480 7.7 案例研究 481 7.7.1 蒙特卡洛积分 481 7.7.2 DNA序列比对 485 第8章 负载均衡 493 8.1 引言 493 8.2 动态负载均衡：Linda的遗赠 494 8.3 静态负载均衡：可分负载理论方法 495 8.3.1 建模开销 496 8.3.2 通信设置 502 8.3.3 分析 503 8.3.4 总结：简短的文献综述 510 8.4 DLTlib：分割工作负载的库 513 8.5 案例研究 516 8.5.1 Mandelbrot集“电影”的混合计算：动态负载均衡案例研究 516 8.5.2 分布式块加密：静态负载均衡案例研究 526 在线资源 附录A 编译Qt程序 附录B 运行MPI程序：准备与配置步骤 附录C 测量时间 附录D Boost.MPI 附录E CUDA环境搭建 附录F DLTlib 术语表 参考文献 ---------------------------3769261 - CUDA并行程序设计：GPU编程指南--------------------------- 《CUDA并行程序设计：GPU编程指南》 致中国读者 译者序 前 言 第1章 超级计算简史 1 1.1 简介 1 1.2 冯·诺依曼计算机架构 2 1.3 克雷 4 1.4 连接机 5 1.5 Cell处理器 6 1.6 多点计算 8 1.7 早期的GPGPU编程 10 1.8 单核解决方案的消亡 11 1.9 英伟达和CUDA 12 1.10 GPU硬件 13 1.11 CUDA的替代选择 15 1.11.1 OpenCL 15 1.11.2 DirectCompute 16 1.11.3 CPU的替代选择 16 1.11.4 编译指令和库 17 1.12 本章小结 18 第2章 使用GPU理解并行计算 19 2.1 简介 19 2.2 传统的串行代码 19 2.3 串行/并行问题 21 2.4 并发性 22 2.5 并行处理的类型 25 2.5.1 基于任务的并行处理 25 2.5.2 基于数据的并行处理 27 2.6 弗林分类法 29 2.7 常用的并行模式 30 2.7.1 基于循环的模式 30 2.7.2 派生/汇集模式 31 2.7.3 分条/分块 33 2.7.4 分而治之 34 2.8 本章小结 34 第3章 CUDA硬件概述 35 3.1 PC架构 35 3.2 GPU硬件结构 39 3.3 CPU与GPU 41 3.4 GPU计算能力 42 3.4.1 计算能力1.0 42 3.4.2 计算能力1.1 43 3.4.3 计算能力1.2 44 3.4.4 计算能力1.3 44 3.4.5 计算能力2.0 44 3.4.6 计算能力2.1 46 第4章 CUDA环境搭建 48 4.1 简介 48 4.2 在Windows下安装软件开发工具包 48 4.3 Visual Studio 49 4.3.1 工程 49 4.3.2 64位用户 49 4.3.3 创建工程 51 4.4 Linux 52 4.5 Mac 55 4.6 安装调试器 56 4.7 编译模型 58 4.8 错误处理 59 4.9 本章小结 60 第5章 线程网格、线程块以及线程 61 5.1 简介 61 5.2 线程 61 5.2.1 问题分解 62 5.2.2 CPU与GPU的不同 63 5.2.3 任务执行模式 64 5.2.4 GPU线程 64 5.2.5 硬件初窥 66 5.2.6 CUDA内核 69 5.3 线程块 70 5.4 线程网格 74 5.4.1 跨幅与偏移 76 5.4.2 X与Y方向的线程索引 77 5.5 线程束 83 5.5.1 分支 83 5.5.2 GPU的利用率 85 5.6 线程块的调度 88 5.7 一个实例——统计直方图 89 5.8 本章小结 96 第6章 CUDA内存处理 99 6.1 简介 99 6.2 高速缓存 100 6.3 寄存器的用法 103 6.4 共享内存 112 6.4.1 使用共享内存排序 113 6.4.2 基数排序 117 6.4.3 合并列表 123 6.4.4 并行合并 128 6.4.5 并行归约 131 6.4.6 混合算法 134 6.4.7 不同GPU上的共享内存 138 6.4.8 共享内存小结 139 6.5 常量内存 140 6.5.1 常量内存高速缓存 140 6.5.2 常量内存广播机制 142 6.5.3 运行时进行常量内存更新 152 6.6 全局内存 157 6.6.1 记分牌 165 6.6.2 全局内存排序 165 6.6.3 样本排序 168 6.7 纹理内存 188 6.7.1 纹理缓存 188 6.7.2 基于硬件的内存获取操作 189 6.7.3 使用纹理的限制 190 6.8 本章小结 190 第7章 CUDA实践之道 191 7.1 简介 191 7.2 串行编码与并行编码 191 7.2.1 CPU与GPU的设计目标 191 7.2.2 CPU与GPU上的最佳算法对比 194 7.3 数据集处理 197 7.4 性能分析 206 7.5 一个使用AES的示例 218 7.5.1 算法 219 7.5.2 AES的串行实现 223 7.5.3 初始内核函数 224 7.5.4 内核函数性能 229 7.5.5 传输性能 233 7.5.6 单个执行流版本 234 7.5.7 如何与CPU比较 235 7.5.8 考虑在其他GPU上运行 244 7.5.9 使用多个流 248 7.5.10 AES总结 249 7.6 本章小结 249 第8章 多CPU和多GPU解决方案 252 8.1 简介 252 8.2 局部性 252 8.3 多CPU系统 252 8.4 多GPU系统 253 8.5 多GPU算法 254 8.6 按需选用GPU 255 8.7 单节点系统 258 8.8 流 259 8.9 多节点系统 273 8.10 本章小结 284 第9章 应用程序性能优化 286 9.1 策略1：并行/串行在GPU/CPU上的问题分解 286 9.1.1 分析问题 286 9.1.2 时间 286 9.1.3 问题分解 288 9.1.4 依赖性 289 9.1.5 数据集大小 292 9.1.6 分辨率 293 9.1.7 识别瓶颈 294 9.1.8 CPU和GPU的任务分组 297 9.1.9 本节小结 299 9.2 策略2：内存因素 299 9.2.1 内存带宽 299 9.2.2 限制的来源 300 9.2.3 内存组织 302 9.2.4 内存访问以计算比率 303 9.2.5 循环融合和内核融合 308 9.2.6 共享内存和高速缓存的使用 309 9.2.7 本节小结 311 9.3 策略3：传输 311 9.3.1 锁页内存 311 9.3.2 零复制内存 315 9.3.3 带宽限制 322 9.3.4 GPU计时 327 9.3.5 重叠GPU传输 330 9.3.6 本节小结 334 9.4 策略4：线程使用、计算和分支 335 9.4.1 线程内存模式 335 9.4.2 非活动线程 337 9.4.3 算术运算密度 338 9.4.4 一些常见的编译器优化 342 9.4.5 分支 347 9.4.6 理解底层汇编代码 351 9.4.7 寄存器的使用 355 9.4.8 本节小结 357 9.5 策略5：算法 357 9.5.1 排序 358 9.5.2 归约 363 9.5.3 本节小结 384 9.6 策略6：资源竞争 384 9.6.1 识别瓶颈 384 9.6.2 解析瓶颈 396 9.6.3 本节小结 403 9.7 策略7：自调优应用程序 403 9.7.1 识别硬件 404 9.7.2 设备的利用 405 9.7.3 性能采样 407 9.7.4 本节小结 407 9.8 本章小结 408 第10章 函数库和SDK 410 10.1 简介 410 10.2 函数库 410 10.2.1 函数库通用规范 411 10.2.2 NPP 411 10.2.3 Thrust 419 10.2.4 CuRAND 434 10.2.5 CuBLAS库 438 10.3 CUDA运算SDK 442 10.3.1 设备查询 443 10.3.2 带宽测试 445 10.3.3 SimpleP2P 446 10.3.4 asyncAPI和cudaOpenMP 448 10.3.5 对齐类型 455 10.4 基于指令的编程 457 10.5 编写自己的内核 464 10.6 本章小结 466 第11章 规划GPU硬件系统 467 11.1 简介 467 11.2 CPU处理器 469 11.3 GPU设备 470 11.3.1 大容量内存的支持 471 11.3.2 ECC内存的支持 471 11.3.3 Tesla计算集群驱动程序 471 11.3.4 更高双精度数学运算 472 11.3.5 大内存总线带宽 472 11.3.6 系统管理中断 472 11.3.7 状态指示灯 472 11.4 PCI-E总线 472 11.5 GeForce板卡 473 11.6 CPU内存 474 11.7 风冷 475 11.8 液冷 477 11.9 机箱与主板 479 11.10 大容量存储 481 11.10.1 主板上的输入/输出接口 481 11.10.2 专用RAID控制器 481 11.10.3 HDSL 483 11.10.4 大容量存储需求 483 11.10.5 联网 483 11.11 电源选择 484 11.12 操作系统 487 11.12.1 Windows 487 11.12.2 Linux 488 11.13 本章小结 488 第12章 常见问题、原因及解决方案 489 12.1 简介 489 12.2 CUDA指令错误 489 12.2.1 CUDA错误处理 489 12.2.2 内核启动和边界检查 490 12.2.3 无效的设备操作 491 12.2.4 volatile限定符 492 12.2.5 计算能力依赖函数 494 12.2.6 设备函数、全局函数和主机函数 495 12.2.7 内核中的流 496 12.3 并行编程问题 497 12.3.1 竞争冒险 497 12.3.2 同步 498 12.3.3 原子操作 502 12.4 算法问题 504 12.4.1 对比测试 504 12.4.2 内存泄漏 506 12.4.3 耗时的内核程序 506 12.5 查找并避免错误 507 12.5.1 你的GPU程序有多少错误 507 12.5.2 分而治之 508 12.5.3 断言和防御型编程 509 12.5.4 调试级别和打印 511 12.5.5 版本控制 514 12.6 为未来的GPU进行开发 515 12.6.1 开普勒架构 515 12.6.2 思考 518 12.7 后续学习资源 519 12.7.1 介绍 519 12.7.2 在线课程 519 12.7.3 教学课程 520 12.7.4 书籍 521 12.7.5 英伟达CUDA资格认证 521 12.8 本章小结 522 |