技术

下一个平台Agent 激发LLM涌现——提示工程 LLM微调理论及实践 大佬沉思 LLM外挂知识库 LLMOps 多模态LLM Python一些比较有意思的库 LLM部分技术源码学习 LangChain源码学习 通用分布式计算引擎Ray Python并发 go依赖注入 go collection gc的基本原理 golang性能分析及优化 数据湖 高性能计算与存储 Linux2.1.13网络源代码学习 《大数据经典论文解读》 三驾马车学习 Spark 内存管理及调优 Yarn学习 从Spark部署模式开始讲源码分析 容器狂占内存资源怎么办? 多角度理解一致性 golang io使用及优化模式 Flink学习 c++学习 学习ebpf go设计哲学 ceph学习 学习mesh kvm虚拟化 学习MQ go编译器以及defer实现 学习go 为什么要有堆栈 汇编语言 计算机组成原理 运行时和库 Prometheus client mysql 事务 mysql 事务的隔离级别 mysql 索引 坏味道 学习分布式 学习网络 学习Linux go堆内存分配 golang 系统调用与阻塞处理 Goroutine 调度过程 重新认识cpu mosn有的没的 负载均衡泛谈 单元测试的新解读 《Redis核心技术与实现》笔记 《Prometheus监控实战》笔记 Prometheus 告警学习 calico源码分析 对容器云平台的理解 Prometheus 源码分析 并发的成本 基础设施优化 hashicorp raft源码学习 docker 架构 mosn细节 与微服务框架整合 Java动态代理 编程范式 并发通信模型 《网络是怎样连接的》笔记 go channel codereview gc分析 jvm 线程实现 go打包机制 go interface及反射 如何学习Kubernetes 《编译原理之美》笔记——后端部分 《编译原理之美》笔记——前端部分 Pilot MCP协议分析 go gc 内存管理玩法汇总 软件机制 istio流量管理 Pilot源码分析 golang io 学习Spring mosn源码浅析 MOSN简介 《datacenter as a computer》笔记 学习JVM Tomcat源码分析 Linux可观测性 学习存储 学计算 Gotty源码分析 kubernetes operator kaggle泰坦尼克问题实践 kubernetes扩缩容 神经网络模型优化 直觉上理解深度学习 如何学习机器学习 TIDB源码分析 什么是云原生 Alibaba Java诊断工具Arthas TIDB存储——TIKV 《Apache Kafka源码分析》——简介 netty中的线程池 guava cache 源码分析 Springboot 启动过程分析 Spring 创建Bean的年代变迁 Linux内存管理 自定义CNI IPAM 共识算法 spring redis 源码分析 kafka实践 spring kafka 源码分析 Linux进程调度 让kafka支持优先级队列 Codis源码分析 Redis源码分析 C语言学习 《趣谈Linux操作系统》笔记 docker和k8s安全访问机制 jvm crash分析 Prometheus 学习 Kubernetes监控 容器日志采集 Kubernetes 控制器模型 容器狂占资源怎么办? Kubernetes资源调度——scheduler 时序性数据库介绍及对比 influxdb入门 maven的基本概念 《Apache Kafka源码分析》——server Kubernetes类型系统 源码分析体会 《数据结构与算法之美》——算法新解 Kubernetes源码分析——controller mananger Kubernetes源码分析——apiserver Kubernetes源码分析——kubelet Kubernetes介绍 ansible学习 Kubernetes源码分析——从kubectl开始 jib源码分析之Step实现 jib源码分析之细节 线程排队 跨主机容器通信 jib源码分析及应用 为容器选择一个合适的entrypoint kubernetes yaml配置 《持续交付36讲》笔记 mybatis学习 程序猿应该知道的 无锁数据结构和算法 CNI——容器网络是如何打通的 为什么很多业务程序猿觉得数据结构和算法没用? 串一串一致性协议 当我在说PaaS时,我在说什么 《数据结构与算法之美》——数据结构笔记 PouchContainer技术分享体会 harbor学习 用groovy 来动态化你的代码 精简代码的利器——lombok 学习 《深入剖析kubernetes》笔记 编程语言那些事儿 rxjava3——背压 rxjava2——线程切换 spring cloud 初识 《深入拆解java 虚拟机》笔记 《how tomcat works》笔记 hystrix 学习 rxjava1——概念 Redis 学习 TIDB 学习 如何分发计算 Storm 学习 AQS1——论文学习 Unsafe Spark Stream 学习 linux vfs轮廓 《自己动手写docker》笔记 java8 实践 中本聪比特币白皮书 细读 区块链泛谈 比特币 大杂烩 总纲——如何学习分布式系统 hbase 泛谈 forkjoin 泛谈 看不见摸不着的cdn是啥 《jdk8 in action》笔记 程序猿视角看网络 bgp初识 calico学习 AQS——粗略的代码分析 我们能用反射做什么 web 跨域问题 《clean code》笔记 《Elasticsearch权威指南》笔记 mockito简介及源码分析 2017软件开发小结—— 从做功能到做系统 《Apache Kafka源码分析》——clients dns隐藏的一个坑 《mysql技术内幕》笔记 log4j学习 为什么netty比较难懂? 递归、回溯、动态规划 apollo client源码分析及看待面向对象设计 学习并发 docker运行java项目的常见问题 OpenTSDB 入门 spring事务小结 分布式事务 javascript应用在哪里 《netty in action》读书笔记 netty对http2协议的解析 ssl证书是什么东西 http那些事 苹果APNs推送框架pushy apple 推送那些事儿 编写java框架的几大利器 java内存模型和jvm内存布局 java exception Linux IO学习 netty内存管理 测试环境docker化实践 netty在框架中的使用套路 Nginx简单使用 《Linux内核设计的艺术》小结 Go并发机制及语言层工具 Linux网络源代码学习——数据包的发送与接收 《docker源码分析》小结 docker namespace和cgroup zookeeper三重奏 数据库的一些知识 Spark 泛谈 链式处理的那些套路 netty回顾 Thrift基本原理与实践(二) Thrift基本原理与实践(一) 回调 异步执行抽象——Executor与Future Docker0.1.0源码分析 java gc Jedis源码分析 深度学习泛谈 Linux网络命令操作 JTA与TCC 换个角度看待设计模式 Scala初识 向Hadoop学习NIO的使用 以新的角度看数据结构 并发控制相关的硬件与内核支持 systemd 简介 quartz 源码分析 基于docker搭建测试环境(二) spring aop 实现原理简述 自己动手写spring(八) 支持AOP 自己动手写spring(七) 类结构设计调整 分析log日志 自己动手写spring(六) 支持FactoryBean 自己动手写spring(九) 总结 自己动手写spring(五) bean的生命周期管理 自己动手写spring(四) 整合xml与注解方式 自己动手写spring(三) 支持注解方式 自己动手写spring(二) 创建一个bean工厂 自己动手写spring(一) 使用digester varnish 简单使用 关于docker image的那点事儿 基于docker搭建测试环境 分布式配置系统 JVM执行 git maven/ant/gradle/make使用 再看tcp kv系统 java nio的多线程扩展 《Concurrency Models》笔记 回头看Spring IOC IntelliJ IDEA使用 Java泛型 vagrant 使用 Go常用的一些库 Python初学 Goroutine 调度模型 虚拟网络 《程序员的自我修养》小结 Kubernetes存储 访问Kubernetes上的Service Kubernetes副本管理 Kubernetes pod 组件 Go基础 JVM类加载 硬币和扑克牌问题 LRU实现 virtualbox 使用 ThreadLocal小结 docker快速入门

架构

大模型推理服务框架 模型服务化(未完成) 大模型RHLF 大模型训练 大模型推理 从Attention到Transformer k8s设备管理 LLM工具栈 ddd从理念到代码 如何应用LLM 小鼠如何驾驭大象(LLM)? 多类型负载协调员Koordinator controller-runtime细节分析 finops学习 kubevela多集群 kubevela中cue的应用 基于k8s的工作流 容器和CPU那些事儿 kubevela源码分析 数据集管理fluid 应用管理平台kubevela karmada支持crd 多集群管理 AutoML和AutoDL 特征平台 实时训练 分布式链路追踪 helm tensorflow原理——python层分析 如何学习tensorflow 数据并行——allreduce 数据并行——ps 机器学习中的python调用c 机器学习训练框架概述 embedding的原理及实践 tensornet源码分析 大模型训练和推理 X的生成——特征工程 tvm tensorflow原理——core层分析 模型演变 《深度学习推荐系统实战》笔记 keras 和 Estimator tensorflow分布式训练 分布式训练的一些问题 基于Volcano的弹性训练 图神经网络 pytorch弹性分布式训练 从混部到统一调度 从RNN到Attention pytorch分布式训练 CNN 《动手学深度学习》笔记 pytorch与线性回归 多活 volcano特性源码分析 推理服务 kubebuilder 学习 mpi 学习pytorch client-go学习 tensorflow学习 提高gpu 利用率 GPU与容器的结合 GPU入门 AI云平台梳理 tf-operator源码分析 k8s批处理调度/Job调度 喜马拉雅容器化实践 Kubernetes 实践 学习rpc BFF openkruise学习 可观察性和监控系统 基于Kubernetes选主及应用 《许式伟的架构课》笔记 Admission Controller 与 Admission Webhook 发布平台系统设计 k8s水平扩缩容 Scheduler如何给Node打分 Scheduler扩展 深入controller openkruise cloneset学习 controller-runtime源码分析 pv与pvc实现 csi学习 client-go informer源码分析 kubelet 组件分析 调度实践 Pod是如何被创建出来的? 《软件设计之美》笔记 mecha 架构学习 Kubernetes events学习及应用 CRI 资源调度泛谈 业务系统设计原则 grpc学习 元编程 以应用为中心 istio学习 下一代微服务Service Mesh 《实现领域驱动设计》笔记 概率论 serverless 泛谈 《架构整洁之道》笔记 处理复杂性 那些年追过的并发 服务器端编程 网络通信协议 架构大杂烩 如何学习架构 《反应式设计模式》笔记 项目的演化特点 反应式架构摸索 函数式编程的设计模式 服务化 ddd反模式——CRUD的败笔 研发效能平台 重新看面向对象设计 业务系统设计的一些体会 函数式编程 《左耳听风》笔记 业务程序猿眼中的微服务管理 DDD实践——CQRS 项目隔离——案例研究 《编程的本质》笔记 系统故障排查汇总及教训 平台支持类系统的几个点 代码腾挪的艺术 abtest 系统设计汇总 《从0开始学架构》笔记 初级权限系统设计 领域驱动理念 现有上传协议分析 移动网络下的文件上传要注意的几个问题 推送系统的几个基本问题 做配置中心要想好的几个基本问题 不同层面的异步 分层那些事儿 性能问题分析 用户认证问题 资源的分配与回收——池 消息/任务队列

标签

k8s设备管理 多类型负载协调员Koordinator controller-runtime细节分析 finops学习 kubevela多集群 kubevela中cue的应用 基于k8s的工作流 容器和CPU那些事儿 kubevela源码分析 数据集管理fluid 应用管理平台kubevela karmada支持crd 多集群管理 helm 从混部到统一调度 volcano特性源码分析 kubebuilder 学习 client-go学习 tf-operator源码分析 k8s批处理调度/Job调度 喜马拉雅容器化实践 Kubernetes 实践 openkruise学习 基于Kubernetes选主及应用 Admission Controller 与 Admission Webhook k8s水平扩缩容 Scheduler如何给Node打分 Scheduler扩展 深入controller openkruise cloneset学习 controller-runtime源码分析 pv与pvc实现 csi学习 client-go informer源码分析 kubelet 组件分析 调度实践 Pod是如何被创建出来的? Kubernetes events学习及应用 CRI 资源调度泛谈 如何学习Kubernetes 以应用为中心 kubernetes operator kubernetes扩缩容 serverless 泛谈 什么是云原生 自定义CNI IPAM docker和k8s安全访问机制 Kubernetes监控 Kubernetes 控制器模型 Kubernetes资源调度——scheduler Kubernetes类型系统 Kubernetes源码分析——controller mananger Kubernetes源码分析——apiserver Kubernetes源码分析——kubelet Kubernetes介绍 Kubernetes源码分析——从kubectl开始 kubernetes yaml配置 CNI——容器网络是如何打通的 当我在说PaaS时,我在说什么 《深入剖析kubernetes》笔记 Kubernetes存储 访问Kubernetes上的Service Kubernetes副本管理 Kubernetes pod 组件

GPU入门

2021年08月18日

简介

CPU 和 GPU - 异构计算的演进与发展世界上大多数事物的发展规律是相似的,在最开始往往都会出现相对通用的方案解决绝大多数的问题,随后会出现为某一场景专门设计的解决方案,这些解决方案不能解决通用的问题,但是在某些具体的领域会有极其出色的表现。

GPU 的架构;内存管理;任务管理;数据类型。

GPU

各种游戏里面的人物的脸,并不是那个相机或者摄像头拍出来的,而是通过多边形建模(Polygon Modeling)创建出来的。而实际这些人物在画面里面的移动、动作,乃至根据光线发生的变化,都是通过计算机根据图形学的各种计算,实时渲染出来的。

图像进行实时渲染的过程,可以被分解成下面这样 5 个步骤:

  1. 顶点处理(Vertex Processing)。构成多边形建模的每一个多边形呢,都有多个顶点(Vertex)。这些顶点都有一个在三维空间里的坐标。但是我们的屏幕是二维的,所以在确定当前视角的时候,我们需要把这些顶点在三维空间里面的位置,转化到屏幕这个二维空间里面。这个转换的操作,就被叫作顶点处理。这样的转化都是通过线性代数的计算来进行的。可以想见,我们的建模越精细,需要转换的顶点数量就越多,计算量就越大。而且,这里面每一个顶点位置的转换,互相之间没有依赖,是可以并行独立计算的
  2. 图元处理。把顶点处理完成之后的各个顶点连起来,变成多边形。其实转化后的顶点,仍然是在一个三维空间里,只是第三维的 Z 轴,是正对屏幕的“深度”。所以我们针对这些多边形,需要做一个操作,叫剔除和裁剪(Cull and Clip),也就是把不在屏幕里面,或者一部分不在屏幕里面的内容给去掉,减少接下来流程的工作量。
  3. 栅格化。我们的屏幕分辨率是有限的。它一般是通过一个个“像素(Pixel)”来显示出内容的。对于做完图元处理的多边形,把它们转换成屏幕里面的一个个像素点。每一个图元都可以并行独立地栅格化
  4. 片段处理。在栅格化变成了像素点之后,我们的图还是“黑白”的。我们还需要计算每一个像素的颜色、透明度等信息,给像素点上色。
  5. 像素操作。把不同的多边形的像素点“混合(Blending)”到一起。可能前面的多边形可能是半透明的,那么前后的颜色就要混合在一起变成一个新的颜色;或者前面的多边形遮挡住了后面的多边形,那么我们只要显示前面多边形的颜色就好了。最终,输出到显示设备。

经过这完整的 5 个步骤之后,完成了从三维空间里的数据的渲染,变成屏幕上你可以看到的 3D 动画了。称之为图形流水线(Graphic Pipeline)。这个过程包含大量的矩阵计算,刚好利用了GPU的并行性。

现代 CPU 里的晶体管变得越来越多,越来越复杂,其实已经不是用来实现“计算”这个核心功能,而是拿来实现处理乱序执行、进行分支预测,以及高速缓存部分。而在 GPU 里,这些电路就显得有点多余了,GPU 的整个处理过程是一个流式处理(Stream Processing)的过程。因为没有那么多分支条件,或者复杂的依赖关系,我们可以把 GPU 里这些对应的电路都可以去掉,做一次小小的瘦身,只留下取指令、指令译码、ALU 以及执行这些计算需要的寄存器和缓存就好了。

于是,我们就可以在一个 GPU 里面,塞很多个这样并行的 GPU 电路来实现计算,就好像 CPU 里面的多核 CPU 一样。和 CPU 不同的是,我们不需要单独去实现什么多线程的计算。因为 GPU 的运算是天然并行的。无论是对多边形里的顶点进行处理,还是屏幕里面的每一个像素进行处理,每个点的计算都是独立的。

一方面,GPU 是一个可以进行“通用计算”的框架,我们可以通过编程,在 GPU 上实现不同的算法。另一方面,现在的深度学习计算,都是超大的向量和矩阵,海量的训练样本的计算。整个计算过程中,没有复杂的逻辑和分支,非常适合 GPU 这样并行、计算能力强的架构。

为什么深度学习需要使用GPU

为什么深度学习需要使用GPU?相比cpu,gpu

  1. gpu核心很多,上千个。
  2. gpu内存带宽更高,速度快就贵,所以显存容量一般不大。因为 CPU 首先得取得数据, 才能进行运算, 所以很多时候,限制我们程序运行速度的并非是 CPU 核的处理速度, 而是数据访问的速度。
  3. 控制流,cpu 控制流很强,alu 只占cpu的一小部分。gpu 则要少用控制语句。现代 CPU 里的晶体管变得越来越多,越来越复杂,其实已经不是用来实现“计算”这个核心功能,而是拿来实现处理乱序执行、进行分支预测,以及高速缓存。GPU 专门用于高度并行计算,因此设计时更多的晶体管用于数据处理,而不是数据缓存和流量控制。GPU 只有 取指令、指令译码、ALU 以及执行这些计算需要的寄存器和缓存。PS: 将更多晶体管用于数据处理,例如浮点计算,有利于高度并行计算。
  4. 编程,cpu 是各种编程语言,编译器成熟。

如图所示,CPU在芯片领域中主要用于降低指令时延的功能,例如大型缓存、较少的算术逻辑单元(ALU)和更多的控制单元。与此相比,GPU则利用大量的ALU来最大化计算能力和吞吐量,只使用极小的芯片面积用于缓存和控制单元,这些元件主要用于减少CPU时延。

CPU / GPU原理与 CUDAGPU 一开始是没有“可编程”能力的,程序员们只能够通过配置来设计需要用到的图形渲染效果(图形加速卡)。在游戏领域, 3D 人物的建模都是用一个个小三角形拼接上的, 而不是以像素的形式, 对多个小三角形的操作, 能使人物做出多种多样的动作, 而 GPU 在此处就是用来计算三角形平移, 旋转之后的位置。为了提高游戏的分辨率, 程序会将每个小三角形细分为更小的三角形,每个小三角形包含两个属性, 它的位置和它的纹理。在游戏领域应用的 GPU 与科学计算领域的 GPU 使用的不同是, 当通过 CUDA 调用 GPU 来进行科学计算的时候, 计算结果需要返回给 CPU, 但是如果用 GPU 用作玩游戏的话, GPU 的计算结果直接输出到显示器上, 也就不需要再返回到 CPU。

深度学习的模型训练,指的是利用数据通过计算梯度下降的方式迭代地去优化神经网络的参数,最终输出网络模型的过程。在这个过程中,通常在迭代计算的环节,会借助 GPU 进行计算的加速。

GPU 架构

和CPU对比

CPU和GPU的主要区别在于它们的设计目标。CPU的设计初衷是执行顺序指令,一直以来,为提高顺序执行性能,CPU设计中引入了许多功能。其重点在于减少指令执行时延,使CPU能够尽可能快地执行一系列指令。这些功能包括指令流水线、乱序执行、预测执行和多级缓存等(此处仅列举部分)。而GPU则专为大规模并行和高吞吐量而设计,但这种设计导致了中等至高程度的指令时延。这一设计方向受其在视频游戏、图形处理、数值计算以及现如今的深度学习中的广泛应用所影响,所有这些应用都需要以极高的速度执行大量线性代数和数值计算,因此人们倾注了大量精力以提升这些设备的吞吐量。我们来思考一个具体的例子:由于指令时延较低,CPU在执行两个数字相加的操作时比GPU更快。在按顺序执行多个这样的计算时,CPU能够比GPU更快地完成。然而,当需要进行数百万甚至数十亿次这样的计算时,由于GPU具有强大的大规模并行能力,它将比CPU更快地完成这些计算任务。

计算和内存架构

理解GPU的底层架构

  1. GPU的core不能做任何类似out-of-order exectutions那样复杂的事情,总的来说,GPU的core只能做一些最简单的浮点运算,例如 multiply-add(MAD)或者 fused multiply-add(FMA)指令,后来经过发展又增加了一些复杂运算,例如tensor张量(tensor core)或者光线追踪(ray tracing core)相关的操作。
  2. GPU的编程方式是SIMD(Single Instruction Multiple Data)意味着所有Core的计算操作完全是在相同的时间内进行的,但是输入的数据有所不同。如果这个工作给到CPU来做,需要N的时间才可以做完,但是给到GPU只需要一个时钟周期就可以完成。
  3. 多个core之间通讯:在图像缩放的例子中,core与core之间不需要任何协作,因为他们的任务是完全独立的。然而,GPU解决的问题不一定这么简单,假设一个长度为8的数组,在第一步中完全可以并行执行两个元素和两个元素的求和,从而同时获得四个元素,两两相加的结果,以此类推,通过并行的方式加速数组求和的运算速度。如果是长度为8的数组两两并行求和计算,那么只需要三次就可以计算出结果。如果是顺序计算需要8次。如果GPU想要完成上述的推理计算过程,显然,多个core之间要可以共享一段内存空间以此来完成数据之间的交互,需要多个core可以在共享的内存空间中完成读/写的操作。我们希望每个Cores都有交互数据的能力,但是不幸的是,一个GPU里面可以包含数以千计的core,如果使得这些core都可以访问共享的内存段是非常困难和昂贵的出于成本的考虑,折中的解决方案是将各类GPU的core分类为多个组,形成多个流处理器(Streaming Multiprocessors )或者简称为SMs。
  4. SM块的底部有一个96KB的L1 Cache/SRAM。每个SM都有自己的L1缓存,SM间不能互相访问彼此的L1。L1 CACHE拥有两个功能,一个是用于SM上Core之间相互共享内存(寄存器 也可以),另一个则是普通的cache功能。存在全局的内存GMEM,但是访问较慢,Cores当需要访问GMEM的时候会首先访问L1,L2如果都miss了,那么才会花费大代价到GMEM中寻找数据。

GPU架构总体如下图所示:

流式多处理器(Streaming Multiprocessor、SM)是 GPU 的基本单元,每个 GPU 都由一组 SM 构成,SM 中最重要的结构就是计算核心 Core

  1. 线程调度器(Warp Scheduler):线程束(Warp)是最基本的单元,每个线程束中包含 32 个并行的线程,GPU 控制部件面积比较小,为了节约控制器,一个 Warp 内部的所有 CUDA Core 的 PC(程序计数器)一直是同步的,但是访存地址是可以不同的,每个核心还可以有自己独立的寄存器组,它们使用不同的数据执行相同的命令,这种执行方式叫做 SIMT(Single Instruction Multi Trhead)。调度器会负责这些线程的调度;
    1. 一个 Warp 中永远都在执行相同的指令,如果分支了怎么处理呢?其实 Warp 中的 CUDA Core 并不是真的永远都执行相同的指令,它还可以不执行。这样会导致 Warp Divergence,极端情况下,每一个Core的指令流都不一样,那么甚至还可能导致一个 Warp 中仅有一个 Core 在工作,效率降低为 1/32.

    2. GPU 需要数据高度对齐,一个 Warp 的内存访问是成组的,一次只能读取连续的且对齐的 128byte(正好是WarpSize 32 * 4 byte),CPU 是一个核心一个 L1,GPU 是两个 Warp 一个 L1 Cache,整个Warp 有一个核心数据没准备好都执行不了。
    3. GPU 的线程切换不同于 CPU,在 CPU 上切换线程需要保存现场,将所有寄存器都存到主存中,GPU 的线程切换只是切换了寄存器组(一个 SM 中有高达 64k 个寄存器),延迟超级低,几乎没有成本。一个 CUDA Core 可以随时在八个线程之间反复横跳,哪个线程数据准备好了就执行哪个。 这是 GPU 优于 CPU 的地方,也是为了掩盖延迟没办法的事情。
  2. 访问存储单元(Load/Store Queues):在核心和内存之间快速传输数据;
  3. 核心(Core):GPU 最基本的处理单元,也被称作流处理器(Streaming Processor),每个核心都可以负责整数和单精度浮点数的计算;
  4. 特殊函数的计算单元(Special Functions Unit、SPU)
  5. 存储和缓存数据的寄存器文件(Register File)
  6. 共享内存(Shared Memory)

与个人电脑上的 GPU 不同,数据中心中的 GPU 往往都会用来执行高性能计算和 AI 模型的训练任务。正是因为社区有了类似的需求,Nvidia 才会在 GPU 中加入张量(标量是0阶张量,向量是一阶张量, 矩阵是二阶张量)核心(Tensor Core)18专门处理相关的任务。张量核心与普通的 CUDA 核心其实有很大的区别,CUDA 核心在每个时钟周期都可以准确的执行一次整数或者浮点数的运算,时钟的速度和核心的数量都会影响整体性能。张量核心通过牺牲一定的精度可以在每个时钟计算执行一次 4 x 4 的矩阵运算。PS:就像ALU 只需要加法器就行了(乘法指令转换为多个加法指令),但为了提高性能,直接做了一个乘法器和加法器并存。

CUDA 编程主打一个多线程 thread,多个 thread 成为一个 thread block,同一个 block 内的 thread 共享Shared Memory/L1 cache/SRAM,而 thread block 就是由这么一个 Streaming Multiprocessor (SM) 来运行的。

  1. 一个 SM 里面有多个 subcore,每个 subcore 有一个 32 thread 的 warp scheduler 和 dispatcher, 在一个 warp 中的所有线程都会同时执行相同的指令,但是输入的数据不同,这种机制也被称为 SIMD(单指令多数据)或 SIMT(单指令多线程)模型。
  2. GPU 的调度单元以 warp 为单位进行调度,而不是单个线程。这意味着整个 warp 会被分配到一个流多处理器(SM)上并一起执行。在 CUDA 中,占用率是一个重要的性能指标,表示每个 SM 上激活的 warps 与 SM 可以支持的最大 warp 数量的比例。更高的占用率通常意味着更好的硬件利用率。
  3. 如果 warp 中的所有线程都采取相同的分支路径(例如,都满足某个条件语句),则它们会继续同步执行。但是,如果线程在分支上有不同的路径(即分歧),则 warp 会执行每个路径,但不是所有线程都会在每个路径上活跃。这可能导致效率下降,因为即使某些线程在特定路径上没有工作,整个 warp 也必须等待该路径完成。为了确保高效执行,开发人员可能需要确保他们的代码减少 warp 分歧。
  4. Global memory 就是我们常说的 显存 (GPU memory),其实是比较慢的。Global memory 和 shared memory 之间是 L2 cache,L2 cache 比 global memory 快。每次 shared memory 要到 global memory 找东西的时候, 会去看看 l2 cache 里面有没有, 有的话就不用去 global memory 了. 有的概率越大, 我们说 memory hit rate 越高, CUDA 编程的一个目的也是要尽可能提高 hit rate. 尤其是能够尽可能多的利用比较快的 SRAM (shared memory).但是因为 SRAM 比较小, 所以基本原则就是: 每次往 SRAM 移动数据的, 都可能多的用这个数据. 避免来来回回的移动数据. 这种 idea 直接促成了最近大火的 FlashAttention. FlashAttention 发现很多操作计算量不大, 但是 latency 很高, 那肯定是不符合上述的 “每次往 SRAM 移动数据的”. 怎么解决呢?Attention 基本上是由 matrix multiplication 和 softmax 构成的. 我们已经知道了 matrix multiplication 是可以分块做的, 所以就剩下 softmax 能不能分块做? softmax 其实也是可以很简单的被分块做的. 所以就有了 FlashAttention.

GPU的线程在执行的时候是分块(block)执行的(通过SIMD执行),所以块(block)内的线程其实是共享pc寄存器,这种执行方式就隐含了,在任何一个时刻,一个块当中的所有线程,会进行一模一样的动作:如果是读内存,那么大家一起读;如果是写内存,那么大家一起写。从而,对于这种情况,相较于各个线程有自己的一片内存区域(按照线程组织数据),按照块组织数据效率会更高。所以我们可以看到,一般给CPU用的数据,都是线性排列的。而给GPU用的数据,基本上都是按照块(对应着GPU编程模型当中的线程块)来组织的。例如,从全局内存中加载的粒度是32*4字节,恰好是32个浮点数,每个线程束中的每个线程恰好一个浮点数。同样的原因,在片上存储空间以及多层高速缓存(cache)的组织方面,GPU也是突出了一个分块交换的概念,对线程组(block)的尺寸是非常敏感的。而这些在CPU上就不是那么明显。GPU深度学习性能的三驾马车:Tensor Core、内存带宽与内存层次结构

Ampere 架构的 GPU一些重要的运算延迟周期时间

  1. 访问全局内存(高达80GB):约380个周期
  2. 二级缓存(L2 cache):约200个周期
  3. 一级缓存或访问共享内存(每个流式多处理器最多128KB):约34个周期
  4. 乘法和加法在指令集层面的结合(fused multiplication and addition,FFMA):4个周期
  5. Tensor Core(张量计算核心)矩阵乘法运算:1个周期
CUDA编程抽象 线程层次结构 存储层次结构
线程 ALU  
线程块 流多处理器 (SM) 共享内存、栅栏同步
网格 GPU 全局内存

Kernel 程序通常以一种拷贝和计算(copy and compute)模式执行,即,首先从全局内存中获取数据,并将数据存储到共享内存中,然后对共享内存数据执行计算,并将结果(如果有)写回全局内存。

执行模型

CPU 与GPU 协作

GPU 无法自己独立工作,其工作任务还是由 CPU 进行触发的。整体的工作流程可以看做是 CPU 将需要执行的计算任务异步的交给 GPU,GPU 拿到任务后,会将 Kernel 调度到相应的 SM 上,而 SM 内部的线程则会按照任务的描述进行执行。

大多数采用的还是分离式结构,AMD 的 APU 采用耦合式结构,目前主要使用在游戏主机中,如 PS4。

  1. 锁页:GPU 可以直接访问 CPU的内存。出于某些显而易见的原因,cpu 和gpu 最擅长访问自己的内存,但gpu 可以通过DMA 来访问cpu 中的锁页内存。锁页是操作系统常用的操作,可以使硬件外设直接访问内存,从而避免过多的复制操作。”被锁定“的页面被os标记为不可被os 换出的,所以设备驱动程序在给这些外设编程时,可以使用页面的物理地址直接访问内存。PS:部分内存的使用权暂时移交给设备。
  2. 命令缓冲区:CPU 通过 CUDA 驱动写入指令,GPU 从缓冲区 读取命令并控制其执行,
  3. CPU 与GPU 同步:cpu 如何跟踪GPU 的进度

对于一般的外设来说,驱动程序提供几个api接口,约定好输入和输出的内存地址,向输入地址写数据,调接口,等中断,从输出地址拿数据。输出数据地址 command_operation(输入数据地址)。gpu 是可以编程的,变成了输出数据地址 command_operation(指令序列,输入数据地址)

系统的三个要素: CPU,内存,设备。CPU 虚拟化由 VT-x/SVM 解决,内存虚拟化由 EPT/NPT 解决,设备虚拟化呢?它的情况要复杂的多,不管是 VirtIO,还是 VT-d,都不能彻底解决设备虚拟化的问题。除了这种完整的系统虚拟化,还有一种也往往被称作「虚拟化」的方式: 从 OS 级别,把一系列的 library 和 process 捆绑在一个环境中,但所有的环境共享同一个 OS Kernel。

不考虑嵌入式平台的话,那么,GPU 首先是一个 PCIe 设备。GPU 的虚拟化,还是要首先从 PCIe 设备虚拟化角度来考虑。一个 PCIe 设备,有什么资源?有什么能力?

  1. 2 种资源: 配置空间;MMIO(Memory-Mapped I/O)
  2. 2 种能力: 中断能力;DMA 能力

一个典型的 GPU 设备的工作流程是:

  1. 应用层调用 GPU 支持的某个 API,如 OpenGL 或 CUDA
  2. OpenGL 或 CUDA 库,通过 UMD (User Mode Driver),提交 workload 到 KMD (Kernel Mode Driver)
  3. Kernel Mode Driver 写 CSR MMIO,把它提交给 GPU 硬件
  4. GPU 硬件开始工作… 完成后,DMA 到内存,发出中断给 CPU
  5. CPU 找到中断处理程序 —— Kernel Mode Driver 此前向 OS Kernel 注册过的 —— 调用它
  6. 中断处理程序找到是哪个 workload 被执行完毕了,…最终驱动唤醒相关的应用

本质上GPU 还是一个外设,有驱动程序(分为用户态和内核态)和API,用户程序 ==> API ==> CPU ==> 驱动程序 ==> GPU ==> 中断 ==> CPU.

warp(gpu的一个单位)是典型的单指令多线程(SIMT,SIMD单指令多数据的升级)的实现,也就是32个线程同时执行的指令是一模一样的,只是线程数据不一样,这样的好处就是一个warp只需要一个套逻辑对指令进行解码和执行就可以了,芯片可以做的更小更快,只所以可以这么做是由于GPU需要处理的任务是天然并行的。

GPU架构与计算入门指南CUDA是NVIDIA提供的编程接口,用于编写运行在其GPU上的程序。

  1. 在CUDA中,你会以类似于C/C++函数的形式来表达想要在GPU上运行的计算,这个函数被称为kernel。
  2. kernel在并行中操作向量形式的数据,这些数字以函数参数的形式提供给它。 一个简单的例子是执行向量加法的kernel,即接受两个向量作为输入,逐元素相加,并将结果写入第三个向量。

要在GPU上执行kernel,我们需要启用多个线程,这些线程总体上被称为一个网格(grid),但网格还具有更多的结构。一个网格由一个或多个线程块(有时简称为块)组成,而每个线程块又由一个或多个线程组成。

线程块和线程的数量取决于数据的大小和我们所需的并行度。例如,在向量相加的示例中,如果我们要对256维的向量进行相加运算,那么可以配置一个包含256个线程的单个线程块,这样每个线程就可以处理向量的一个元素。如果数据更大,GPU上也许没有足够的线程可用,这时我们可能需要每个线程能够处理多个数据点。

编写一个kernel需要两步。第一步是运行在CPU上的主机代码,这部分代码用于加载数据,为GPU分配内存,并使用配置的线程网格启动kernel;第二步是编写在GPU上执行的设备(GPU)代码。在GPU上执行Kernel的步骤

  1. 将数据从主机复制到设备。 在调度执行kernel之前,必须将其所需的全部数据从主机(即CPU)内存复制到GPU的全局内存(即设备内存)。在最新的GPU硬件中,我们还可以使用统一虚拟内存直接从主机内存中读取数据。
  2. SM上线程块的调度。当GPU的内存中拥有全部所需的数据后,它会将线程块分配给SM。同一个块内的所有线程将同时由同一个SM进行处理。为此,GPU必须在开始执行线程之前在SM上为这些线程预留资源。在实际操作中,可以将多个线程块分配给同一个SM以实现并行执行。由于SM的数量有限,而大型kernel可能包含大量线程块,因此并非所有线程块都可以立即分配执行。GPU会维护一个待分配和执行的线程块列表,当有任何一个线程块执行完成时,GPU会从该列表中选择一个线程块执行。
  3. 单指令多线程 (SIMT) 和线程束(Warp)。众所周知,一个块(block)中的所有线程都会被分配到同一个SM上。但在此之后,线程还会进一步划分为大小为32的组(称为warp),并一起分配到一个称为处理块(processing block)的核心集合上进行执行。SM通过获取并向所有线程发出相同的指令,以同时执行warp中的所有线程。然后这些线程将在数据的不同部分,同时执行该指令。在向量相加的示例中,一个warp中的所有线程可能都在执行相加指令,但它们会在向量的不同索引上进行操作。由于多个线程同时执行相同的指令,这种warp的执行模型也称为单指令多线程 (SIMT)。这类似于CPU中的单指令多数据(SIMD)指令。 PS: 有点类似hadoop 任务调度的意思,cpu 是driver 进程,gpu是 worker 进程

CUDA——GPU 编程

NVIDIA 率先在 GPU 中引入了通用计算能力,使得开发者能利用 CUDA 编程语言来驱动。这时候 GPU 的核心都是 CUDA Core。由于一个 GPU 里面有大量的 CUDA Core,使得并行度高的程序获得了极大的并行加速。但是,CUDA Core 在一个时钟周期只能完成一个操作,矩阵乘法操作依然需要耗费大量的时间。NVIDIA 为了进一步加速“加乘运算”,在 2017 年推出了 Volta 架构的 GPU,从这个架构开始 Tensor Core 被引入。它可以在一个时钟周期完成两个 4×4x4 半精度浮点矩阵的乘法(64 GEMM per clock)。

CUDA C++ 编程指北-第一章:入门以及编程模型 未读。

CUDA编程入门极简教程2006年,NVIDIA公司发布了CUDA,CUDA是建立在NVIDIA的CPUs上的一个通用并行计算平台和编程模型。CUDA编程模型是一个异构模型,需要CPU和GPU协同工作。CUDA的架构中引入了主机端(host, cpu)和设备(device, gpu)的概念,我们用host指代CPU及其内存,而用device指代GPU及其内存,CUDA程序中既包含host程序,又包含device程序,它们分别在CPU和GPU上运行。CUDA的Kernel函数既可以运行在主机端,也可以运行在设备端。同时主机端与设备端之间可以进行数据拷贝。

device 函数和global函数因为需要在GPU上运行,因此不能调用常见的一些 C/C++ 函数(因为这些函数没有对应的 GPU 实现)。

限定符 执行 调用 备注
global 设备端执行 可以从主机调用也可以从某些特定设备调用 异步操作,host 将并行计算任务发射到GPU的任务调用单之后,不会等待kernel执行完就执行下一步
device 设备端执行 设备端调用  
host 主机端执行 主机调用  

典型的CUDA程序的执行流程如下:

  1. 分配host内存,并进行数据初始化;
  2. 分配device内存,并从host将数据拷贝到device上;
  3. 调用CUDA的核函数在device上完成指定的运算;
  4. 将device上的运算结果拷贝到host上;
  5. 释放device和host上分配的内存。

矩阵加法示例

// __global__ 表示在device上执行从host中调用
// 两个向量加法kernelgrid和block均为一维
__global__ void add(float* x, float * y, float* z, int n){
    // 获取全局索引
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    // 步长
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride){
        z[i] = x[i] + y[i];
    }
}
int main(){
    int N = 1 << 20;
    int nBytes = N * sizeof(float);
    // 申请host内存
    float *x, *y, *z;
    x = (float*)malloc(nBytes);
    y = (float*)malloc(nBytes);
    z = (float*)malloc(nBytes);
    // 初始化数据
    for (int i = 0; i < N; ++i){
        x[i] = 10.0;
        y[i] = 20.0;
    }
    // 申请device内存
    float *d_x, *d_y, *d_z;
    cudaMalloc((void**)&d_x, nBytes);
    cudaMalloc((void**)&d_y, nBytes);
    cudaMalloc((void**)&d_z, nBytes);
    // 将host数据拷贝到device
    cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);
    // 定义kernel的执行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    // 执行kernel
    add << < gridSize, blockSize >> >(d_x, d_y, d_z, N);
    // 将device得到的结果拷贝到host
    cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);
    // 检查执行结果
    float maxError = 0.0;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(z[i] - 30.0));
    std::cout << "最大误差: " << maxError << std::endl;
    // 释放device内存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    // 释放host内存
    free(x);
    free(y);
    free(z);
    return 0;
}

如何在 CPU 之上调用 GPU 操作?可以通过调用 __global__ 方法来在GPU之上执行并行操作。我的第一份CUDA代码 - xcyuyuyu的文章 - 知乎

CUDA编程基础与Triton模型部署实践 未读。

GPU/CUDA/驱动和机器学习训练框架的关系

显卡是硬件,硬件需要驱动,否则不能调用其计算资源。CUDA又是什么?

  1. 在2007年之前,GPU由CPU操作,CPU把一些图形图像的计算任务交给GPU执行。程序员不需要与GPU打交道。随着GPU计算能力的发展,越来越多的计算场景由GPU完成效果会更好。但现有的程序无法直接自由控制GPU的处理器。当然程序员也可以直接写代码与显卡驱动对接,从而直接控制GPU的处理器,但这样代码恐怕写起来要让人疯掉。nvidia当然会有动力提供一套软件接口来简化操作GPU的处理器。nvidia把这一套软件定义为CUDA。
  2. 多核 CPU 和众核 GPU 的出现意味着主流处理器芯片现在是并行系统。挑战在于开发能够透明地扩展可并行的应用软件,来利用不断增加的处理器内核数量。CUDA 并行编程模型旨在克服这一挑战,同时为熟悉 C 等标准编程语言的程序员保持较低的学习曲线。CUDA 编程手册系列第一章:CUDA 简介

gpu 和 cuda 和 gpu driver 之间的关系:比如 TX3090需要Compute Capability在8.6以上的cuda,而满足这个要求的cuda又只有11.0以上的版本。而cuda11版本又需要版本号>450的显卡驱动。

显卡  
    ==> Compute Capability   查看显卡支持的Compute Capability, https://developer.nvidia.com/cuda-gpus
        ==> cuda 
            ==> GPU driver   查看cuda对驱动的要求 (https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html
            ==> tf/pytorch version

Compute Capability的数值和GPU的计算速度无关,但是和GPU可执行的任务种类有关。The compute capability of a device is represented by a version number, also sometimes called its “SM version”. This version number identifies the features supported by the GPU hardware and is used by applications at runtime to determine which hardware features and/or instructions are available on the present GPU.