AI-Infra HPC 学习与总结

本仓库用于记录 AI-Infra 与 HPC 技术：

• AI System 的底层技术
• 集群多机多卡互联技术
• 并行计算与高性能计算
• 训练与推理

文件夹说明

各文件夹的内容如下：

├── 01 chip		# 芯片硬件
├── 02 hpc		# 高性能计算
│   ├── 01 openmp	
│   ├── 02 simd
│   └── 05 cuda		# GPU 编程
├── 03 link		# 多机多卡互联底层通信
│   ├── 01 noc		# 片上网络
│   ├── 02 pcie		# PCI-Express(peripheral component interconnect express)
│   ├── 03 topo		# 多机多卡互联以及拓扑相关
│   ├── 05 gpuDirect 	# GPU 与 Host、GPU、IB 网卡、NVMe SSD 的底层通信
│   └── 08 infiniband 	# IB 网卡与 RDMA 通信
├── 04 storage		# 存储
├── 05 ccl 		# 集合通信
│   ├── mpi 		# MPI 的使用与教程
│   └── nccl 		# NCCL 的使用、设计
├── 06 trainAndInfer 	# 训练与推理
├── 99 industryReport 	# 行业前沿报告

CUDA 总结

• CUDA 编程
• 1.CUDA中的异构并行计算
  • GPU核心和CPU核心
• 2.CUDA 编程模型
  • 异构架构
  • 内存层次结构
  • 线程管理
    • CUDA为什么要分线程块和线程网格？
  • 函数类型限定符
    • CUDA核函数的限制
  • CUDA中的修饰符
    • __global__
    • __device__
    • __host__
    • __host__ __device__
    • __noinline__
    • __forceinline__
    • __constant__
    • __shared__
    • __managed__
    • __restrict__
• 3.CUDA 执行模型
  • GPU架构概述
    • SM （流多处理器）
    • Block 是分配到 SM 的最小单位
      • Block 无法跨 SM 拆分的原因
    • 线程束
    • SIMD vs SIMT
  • CUDA编程的组件与逻辑
  • 通过 Fermi 架构理解线程执行
  • 线程束执行的本质
    • 线程束和线程块
    • 线程束分化
    • 资源分配对SM中线程束与线程块数量的影响
    • 常驻线程块
• 4.内存布局
  • 内存模型
    • 寄存器（Registers）
      • 寄存器的容量
      • 寄存器的大小
    • 本地内存（Local Memory）
    • 共享内存（Shared Memory）
    • 常量内存（Constant Memory）
    • 纹理内存（Texture Memory）
    • 全局内存（Global Memory）
  • GPU缓存
  • 静态全局内存
  • 固定内存
  • 零拷贝内存
  • 统一虚拟寻址（UVA）
  • 统一内存（UM）
    • UM 与 UVA 的区别
  • 内存访问模式
    • 对齐与合并访问
• 5.共享内存与常量内存
• 6.流和并发
  • 流和事件概述
  • CUDA流
  • 流调度
    • 虚假的依赖关系
    • Hyper-Q技术
  • 流的优先级
  • CUDA事件
    • 事件创建和销毁
    • 记录事件和计算运行时间
  • 流同步
    • 阻塞流和非阻塞流
    • 隐式同步
    • 显式同步
    • 可配置事件
  • 重叠内核执行与数据传输
  • 重叠GPU和CPU执行
  • 流回调
• 其它
• GPU中的内存架构是什么？
  • 1. SP（流处理器）与寄存器文件（Register File）
  • 2. SM（流式多处理器）独享资源
    • (1) L1缓存
    • (2) 共享内存（Shared Memory）
  • 3. 全局共享资源
    • (1) L2缓存
    • (2) 全局内存（Global Memory）
  • 4. 其他内存类型补充
  • 层级关系示意图
  • 总结与修正
• GPU与CPU的同步与内存屏障机制
• Nsight性能分析
• 无图形界面时的性能分析
  • 查看默认的报告
  • 查看指定的报告
  • 报告类型
  • 查看某个报告的详细信息
• CUDA-GDB 调试
• 编译
• 启动调试器
• 基本调试命令
  • 程序控制
  • 断点设置
  • 查看变量和内存
  • 调用栈
• CUDA 特定命令
  • CUDA 线程和块信息
  • CUDA 线程切换
  • CUDA 内核调试
  • CUDA 设备管理
  • CUDA 内存检查
  • CUDA Warp 和 Lane 操作
  • 多 GPU 调试
• 高级技巧
  • 1. 条件断点
  • 2. 观察点（Watchpoint）
  • 3. 命令序列
  • 4. 日志记录
  • 5. 多线程调试
  • 6. 检查内存错误
  • 7. 远程调试
    • 完整步骤
    • 使用 SSH 隧道（推荐）
    • 调试已运行的程序
  • 8. 调试脚本
• 调试场景示例

互联总结

• NVLlink 与 NVSwitch 发展
  • NVLink 发展
  • NVSwitch 发展
• GPU 互联内部结构
• NVLink 技术细节
  • 初代 NVLink 结构
  • 八块 P100 GPU 互联拓扑
• NVSwitch 技术细节
  • 初代 NVSwitch 结构
  • 初代 NVSwitch Block
  • NVSwitch 简化原理与特性
  • 第三代 NVSwitch
• 新华三类脑服务器PCIe拓扑结构说明
  • AI串行配置
  • AI并行配置
  • HPC配置
  • 三种拓扑对比选型
• 海康类脑服务器PCIe拓扑结构说明
  • QPI与UPI

GPUDirect 总结

• GPUDirect
  • GDS（GPUDirect Storage）
  • GPUDirect P2P
  • GDR（GPUDirect RDMA）
    • GPU上的RDMA操作
    • InfiniBand GPUDirect Async
• 设备侧ARM发起RDMA通信
  • 通信资源由谁控制？
  • 将host端内存映射到device侧，使设备侧可以访问主机侧内存
  • arm提交rdma网络请求
    • 发送与接收
    • 轮询完成队列
• GDRCopy
• NVSHMEM
  • NVSHMEM 和 NCCL
• 参考
• 【研究综述】浅谈GPU通信和PCIe P2P DMA
• GPU内存管理
  • 分离内存管理
    • 锁页内存
    • 零拷贝内存
  • 半分离内存管理——统一虚拟地址空间 UVA
  • 统一内存管理——UM
• GPUDirect 技术的演进
• GPUDirect 技术的实现细节（三个关键问题）（核心章节）
  • 关键问题一：网卡如何直接读写GPU显存？(CPU控制通信)（原理一定要懂）
  • 关键问题二：GPU如何访问通信资源（GPU 控制通信）
    • 通信资源的划分
    • GPU如何访问网卡寄存器
    • GPU如何访问队列资源
  • 关键问题三：GPU如何提交通信请求并与网卡同步
• 总结 GPUDirect 技术的实现细节

RDMA 总结

• RDMA 原理
• 基本概念
  • 缩写
  • 什么是 RDMA
    • 硬件结构
    • RDMA 特点
  • RDMA 技术优缺点
  • 支持 RDMA 的网络协议
    • RoCE适配器可以与其他适配器类型（例如iWARP）通信吗？
    • RoCE出现的背景
    • RoCE和IB网络层和链路层协议区别
    • RDMA 、InfiniBand、IBoE、RoCE、iWARP、IB卡、IB驱动的关系
  • RoCEv1与RoCEv2性能的差异
• DPDK和RDMA的区别
  • 基于 OS 内核的数据传输的弊端
  • dpdk 相对传统基于 OS 内核的数据传输的改进
• RDMA的编程模型更贴近硬件的实现
• DoorBell
• 拥塞控制
• SEND & RECV
• WRITE 和 READ
  • WRITE
  • READ
• 立即数的传输方式
• 连接
  • 连接阶段获取的对端信息
  • 基于连接
  • 基于数据报
  • 服务类型
  • 资源消耗
• 内存注册
  • 注册 GPU 内存
• Protection Domain
• Address Handle
• QPC
• CQ & CQC
  • CQC
  • CQ
  • 可靠服务RC的CQE产生过程
  • 不可靠服务UD的CQE产生过程
  • CQ的保序
• 用户态与内核态的交互
• MW（memory windows）
  • 为什么要有MW
  • MW和MR权限的关系
• RDMA 性能优化
  • 小数据优化延迟原理
    • 快速了解IB通信中的PCIe原理
    • 通信机制
  • 小数据优化延迟方法
    • 避免数据路径中出现控制面操作。
    • 一次 post 多个 WR：
    • Inlining/内联（IBV_SEND_INLINE）：
      • 使用INLINE的陷阱
    • Unsigned Completions（不设置 IBV_SEND_SIGNALED）
    • 与 IB 网卡同 NUMA
    • 其它未验证的
  • 优化带宽的方法
    • MTU
  • 减少地址翻译的性能开销
  • 建立多个QP来加速数据处理
  • 性能优化参考
• 带宽问题分析
  • 服务器架构（拓扑结构）对性能的影响
  • NUMA 非一致性内存访问
  • 芯片间的拓扑结构
  • 查看指定设备与cpu和NUMA 节点的亲和性
    • /proc/bus/pci/devices
    • lspci -v
    • lspci | grep 设备代号
• 降低RDMA延迟波动的方法
  • CPU隔离
  • 减少该cpu上的中断
  • 把该cpu从宽限期隔离出来
• RDMA 链式工作请求与门铃机制分析
  • SGE方式的配置限制
  • 链式提交
  • 一次提交的WR超过SQ剩余WR的数量会报错
  • 链式提交只触发一次门铃
• ibv_send_flags
• qp_init_attr.sq_sig_all的作用
• 如果sq_sig_all=0，每隔 N 个请求设置一次IBV_SEND_SIGNALED，发送队列的头尾指针更新过程
• 直接传输网卡设备内存上的数据
• RDMA 环境
• RDMA 环境部署
  • 使用 OFED 包安装 IB 驱动
  • 安装 rdma 依赖包
  • 查看动态库属于哪一个包的工具
• 常用命令
  • 性能测试
  • 查看 ib 端口
  • 用lspci查看ib的带宽等硬件信息
  • 修改速率
• 错误处理
  • 常见错误码
  • PORT DOWN问题
  • PORT INIT 问题
  • 解除锁定的内存最大为 64MB 的限制

NCCL 总结

• 算法
• 算法原理与实现
  • Broadcast Ring
  • Broadcast Tree
  • AllGather Full Mesh
  • AllGather Ring
  • AllGather Bruck
  • Reduce Tree
  • ReduceScatter Ring
  • AllReduce Ring
• 算法耗时总结
• NCCL源码分析
• UniqueId的创建
• bootstrap网络连接的建立
• 机器内拓扑分析
• 建图过程
• 路径计算
• channel 搜索
• 机器间channel连接
• 数据通信链路transport的建立
• 单机内ncclSend和ncclRecv的过程
• 多机间ncclSend和ncclRecv的过程
• double binary tree
• NCCL Protocol
  • Simple协议
  • L(ow)L(atency)协议
  • L(ow)L(atency)128协议
  • Simple 协议
  • LL 和 LL128 协议的对比
  • 为什么 LL128 协议的数据块较大？
• 千卡训练经验
• 其它命令
• 查看拓扑

训练与推理总结

• 卷积
• NLP领域中的token和tokenization到底指的是什么？
  • 什么是Tokenizer？
• 大模型参数量和显存的换算关系
  • 1B参数对应多少G显存？
  • 训练时的显存开销
    • 激活值的数据占了显存的大头
    • 激活值
  • 推理时的显存开销
• MFU(Model FLOPS Utilization）
• 目前训练超大规模语言模型技术路线：GPU + PyTorch + Megatron-LM + DeepSpeed。这几个工具在训练中各自的作用是什么？
• 并行
  • 数据并行
    • DP
      • DP通讯瓶颈与梯度异步更新
    • DDP
  • 模型并行
    • Tensor并行
    • 流水线并行
• ZeRO
• DeepEP
  • 专家并行的通信需求
  • DeepEP的主要特性
  • DeepEP对外接口