CS149ParallelComputing_LectureNotes

视频链接

课程网站

Lecture1 - 10，尽量只记录对Assignment可能有帮助的部分，Assignment1 - 4 需要到的内容基本都在前10节课

Lecture11 以后的部分随便看看

Lecture1 Why parallelism Why Efficiency

在18年前，处理器效率提升的两个重大因素：

  • 指令级并行（instruction level parallelism）
  • 提高CPU频率

指令级并行是什么？这里先回顾单核无指令集并行的处理器模型，三个部分组成：取指\译指器、ALU、ExectionContext(寄存器组合)，这种处理串行地执行代码。

image-20260615230813084

但代码中有许多没有依赖关系的部分，指令级并行的关键就是识别这些没有依赖关系的指令，比如在下面的程序中，x、y、z的计算互不依赖，只是最后需要把前三个乘法运算相加，这一步骤依赖前序三个乘法运算。

image-20260615231105107

可以根据依赖关系并行执行指令：

image-20260615231704476

但实际上我们不需要两个processor，只需要两套取指\译指器和两套ALU，如下图所示，Superscalar proccessor 可以在每个周期执行两个指令
image-20260615231930818

我们可以再增加一个核中取指\译指器、ALU的数量，进一步取得效率提升，但是呈现边际收益递减现象

image-20260615232137350

所以指令级并行的效率提升有瓶颈，再加上”功耗墙“的存在，18年前的效率提升手段逐渐失效了，仅仅硬件升级就能带动软件效率提升的时代渐行渐远了，软件开发者必须考虑如何“并行地”写软件来取得性能提升，而且需要了解各种特殊的加速硬件（gpu、npu等）才能写出运行良好的并行软件。以上就是学习这门课的动机之一。

效率提升几乎总是和”高效获取内存“有关，这也是缓存存在的原因，下图对比了从不同层级内存获取数据的耗时：

image-20260615234237056

Lecture2 Modern Multi-Core Processor

本节课，继续深入讲解现代多核系统的设计。

上一节讲到Superscalar proccessor 在同一个核上集成多个ALU和译码器，达到执行效率提升的目标。此外，还有两种优化手段：

1. Multi-Core

image-20260617232112703

软件如果想要利用这种并行性，需要额外的API -- 比如linux的pthread

2. SIMD（Single Instructtion MultiIple Data）

增加ALU，所有的ALU都将执行相同的指令；各个厂家的处理器都有自己的一套intrinsic datatype和 functions，如IntelX86 的SSE、AVX， ARM的Neon等

image-20260617232753152

涉及SIMD时，尤其要注意if分支引起的divergent execution，在一个支持8-wide SIMD操作的核上，最坏可能达到峰值性能的八分之一

image-20260617234044002

3. Hardware Multi-threading

上一节最后提到，软件性能很大程度上取决于内存获取的速读，直接从Cache中获取速读最快。如果数据能经常Cachehit，将有效减少cpu等待io的时间，减少CPU“卡顿” （stall）；除了Cache，prefetch（预取）数据也能减少CPU“卡顿”（实际没有减少，只是隐藏（hide）了）

但如果数据不在Cache，也不能很好的预测接下来指令流要获取哪些内存呢？ —— 在同一核上交替执行不同线程可以减缓CPU卡顿，如果一个线程的内存迟迟等不来，那就运行另一个线程，这种做法实际上也没有减少io时间，只是被隐藏了。这种做法要求在同一个核上集成多个Context上下文（寄存器），

image-20260618000659301

实现同一个核跑多个线程的方式大概有两种，两种方式都增加了Context数量

  1. interleaved Multi-threading
  2. simultaneous multi-threading， intel的超线程（Hyper-Threading）就是指这个

并行技术总结

第一二讲主要介绍四种并行技术：

  • superscalar，指令集并行
  • SIMD，单指令多数据
  • multi-core，多核
  • hardware multi-threading，硬件层的多线程

这几种技术可以同时出现在一个CPU处理器架构中，现代CPU大部分都是这种模型
image-20260618002416290

对GPU来说，其底层架构是类似的，但CUDA提出了新的编程模型“SIMT”，实际上在硬件层面可能更简单了
image-20260618002527570

Lecture3 Latency and Bandwidth & ISPC

Latency and Bandwidth

MemoryBandwidth(内存带宽宽）：The rate at which the memory system can provide data to a processor.

不要把带宽和延迟混淆，以道路和车进行类比，如下图所示，若车的速读为100km/h，那么带宽是 2/hr, 延迟是0.5h

image-20260621183850681

有三种种方式提升带块，一是提高单车行驶速度，二是增加道路条数：

image-20260621184023903

三是更加效率地利用道路，把车看做内存数据，下图下半部分类似于现代内存传输模型。

image-20260621184102143

内存搬运相对于CPU运算来说，是非常耗时的，所以现代计算机应用的计算效率受带宽限制

image-20260621184327520

当我们说“一个cpu核在一个时钟周期内，执行一条指令”时，我们值得是 “Instruction throughput”而不是“Instruction latency”：

image-20260621184555508

ISPC

Intel SPMD Program Complier （ISPC）

SPMD ： Simple Program Multiple Data

SPMD 编程抽象：对ISPC函数的调用会产生”一系列的程序实例“ （a gang of ISPC program instance，不知道咋翻译）

ISPC语言关键字：

  • programCount：同时运行的ispc程序实例的数量
  • programIndex：当前程序在这个“gang”中的索引
  • uniform：被uniform修饰的变量，表示当前“gang”中的所有程序实例都有相同的值

image-20260621185300178

按照上图右侧的ISPC程序，每个程序实例计算的outputarray位置如下：
image-20260621185848032

不记ISPC了吧，如果只是做PA的话随便看看例子、网络的资料就行，重要的是理解它的编程模型：

  • 一次调用会产生N个程序实例，程序员可以使用几个关键词对程序实例进行区分，底层实现使用SIMD。

Lecture5 Performance Optimization I : work distribution and scheduling

关键目标：

  • 负载均衡
  • 减少通信
  • 减少额外的动作（overhead）以增加并行性，管理作业，减少通信等

Work distribution 策略：

1.静态分配（Static Assignment）

以Assignment1为例，可以均等得将图片从上倒下分割，将每个块分发给不同的线程。

这样做的优点：简单，在作业分发上无额外的计算（overhead）

缺点：极易造成负载不均衡导致并行性差

什么时候适合使用静态分配？ —— 当工作负载可预测时

2.动态分派（Dynamic Assignment）

可以将task分成几个sub_task, 将subt task放在 workqueue 中，各worker thread从workqueue中取task工作

image-20260621213429752

如何划分sub_task的“颗粒度”(granularity)也是有讲究的：

  • 划分的sub_task个数比处理器个数多对整个系统的运行有利
  • 但是更少的sub_task个数减少了管理subtask的overhead
  • 理想的sub_task颗粒度依赖多种因素

为了减少各workthread的同步，也能缓解潜在的负载不均衡问题，可以实现更加“聪明的”动态分配，即每个workthread有一个workqueue，每当一个workthread无事可做时，可以从别的workqueue中“偷”一个task
image-20260621214701881

scheduling的部分略过

Lecture 6 Performance Optimization II ： locality， communication，contension

略

Lecture7 GPU architecture and CUDA Programming

其他更多请参考：CUDA Programming Guide — CUDA Programming Guide

基本概念以及编程模型

基本概念

Cuda编程语言是对GPU硬件的一种“非图形学特定”（non-graphics-specific）的编程接口，在2007年 NVIDIA Tesla 架构中亮相。

Cuda程序由多层级的并发线程构成。线程ID可以至多是3维的（下图例子是2维）

image-20260610000202769

launch 一个cuda kernel函数的写法如下图所示，<<< >>> 中指定gridnum和blocknum。以下图为例，maxtrixAdd <<<numblocks, threadPerBlock>>>(A, B, C) launch的grid维度为3*2， 每个grid包含4 * 3个“线程”，每个线程都会运行函数maxtrixAdd 中的代码，且每个线程在执行过程中都可以用blockIdx、blockDim以及threadIdx这几个内置变量得到自身所处的“位置”。
image-20260610001948803

Cuda内存模型

  1. 宿主机（host）内存空间和设备（gpu）内存空间是分开的。如果要把宿主机上的一份数据传递到gpu中，需要先在gpu中分配一段内存，然后再讲数据拷贝到gpu中

  2. gpu内存也分3种：每个线程自己的内存，每个block的内存（由block中的所有线程共享， 由关键词__shared__标识）以及全局内存（有 cudaMalloc分配）。这三种内存也对应了三种不同的“局部性”（locality），shared内存对性能更友好

    image-20260610003146150

同步：

  1. __syncthreads(): 一种Barrier，等block的所有线程都执行到调用__syncthread()那一行，再往下执行剩余的程序
  2. 原子操作，比如 atomicAdd
  3. Host/Device 同步：kernel函数返回时所有线程之间存在隐式barrier

NVIDIA V100 硬件架构

NVIDIA V100 Gpu 一共有80个 SM（Streaming Multiprocessor），它们共享一个L2Cache。

image-20260610230718015

每个SM都由四个sub-core组成，每个sub_core配有至多16 * 32套执行上下文（R0、R1...寄存器，都是scalar的），每32套上下文组成一个“warp”。如下图所示，每个SM都有如下结构：

image-20260610232202925

每个sub-core有一个Warpselector，运行阶段sub-core选择一个可运行的“warp”进行运算，为该warp的所有线程取得下一个（且同一个）instruction并运算（有些线程可能不运行，这取决于该warp中线程的diverge程度），每个sub-core都配有几个运行单元，如下图所示。

image-20260610230327682

虽然cuda号称使用的是“SIMT”（single instruction multiple thread）编程模型，但如果同一个warp的32个线程都执行同一个指令，事实上就是一种“SIMD”运行方式，且类似于ISPC，执行流divergence也会导致性能下降 —— “If the 32 CUDA threads do not share the same instruction, performance can suffer due to divergent execution”

Lecture 10 Efficiently Evaluating DNNS on Gpus

Lecture 11 Cache Coherence

三种cache miss：

  • Cold Miss -- 冷启动导致的
  • Capacity Miss -- cache容量小导致的
  • Conflict Miss —— cache实现方式导致的（非全相联，not fully associative），比如 8路组相联缓存（可参考这个链接）意味着：一个内存缓存行只能在整个cache中的某8个位置，这提高了查询的效率，但是导致了更多由于冲突引起的cache miss

Cache Design

Write through 和 Write back

  • write through： 写入缓存的同时也写入内存
  • write back：只写缓存，之后再把数据写会内存（需要 dirty 标志位）

write-allocate 和 write no allocate

image-20260622230124727

在一个内存共享的多核处理器上，一个合理的预期是：读一个地址为X的变量，返回最后一次写入X地址的值（最新值）
image-20260622230554853

缓存一致性的挑战：

  1. 使用锁是不能解决缓存一致性问题的，因为每个处理器都有自己的cache，每个cache都可能存储了同一个内存地址的值拷贝
  2. 缓存一致性问题甚至在一个单处理器核上也有可能出现：涉及带有DMA的IO操作时

image-20260622231254050

一些直觉上的问题

“读一个地址为X的变量，返回最后一次写入X地址的值（最新值）”

  • “最后一次”（或者说“最新“），到底指什么？如果两个处理器在同一时间写入怎么办？如果处理器1的写紧跟着处理器2的读，它们的操作在时间上太接近了，以至于不能及时地告诉处理器2读的值被处理器1改过，这样的情况怎么办？
  • “程序的运行顺序（不是时间顺序）决定谁才是“最新”的“ 。 在单线程情况下，这句话是对的，但在多核多线程环境下，需要一种更合理的方式描述“运行顺序”

缓存一致性的定义如下：

image-20260622231957545

实现缓存一致性需要两个不变量（invariants）：

对任意内促地址X，在一个时间段内（epoch），

  • SingleWrite， Multiple-Read（SWMR Invariant）
    • 读写周期（readwrite epoch）：只能有一个处理器能够写入地址X（当然也能读）
    • 只读周期（readonly epoch）：所有处理器只读
  • Data Value Invariant：在某个epoch中，地址X里存储的值与上一个read-write opoch的值相同

image-20260622233104845

实现缓存一致性的具体方式有软件和硬件两种

  • software-based：不讨论
  • hardware-based：
    • “snooping”（嗅探） based （重点提及）
    • directory based

基于”嗅探“的缓存一致性框架

  • 主要思想：所有缓存相关的操作都被广播到系统中的所有处理器的缓存控制器中
  • 缓存控制器监测相关的缓存操作是否与当前处理器相关，并且遵循缓存控制协议（cache coherence protocol）来维持缓存一致性

cache coherencewith write-back cahce

  • 使用总线，总线有两个特点：
    • 同一时间只允许一个transaction，这意味着串行性（间接保证了Data Value Invariant）
    • 能广播所有消息到所有处理器
  • 这要求：dirty标志位只能出现一个缓存中（保证了SWMR Invariant）。如何保证只存在一个缓存中？ ——> 缓存一致性协议（cache coherence protocal）

image-20260623000203981

缓存一致性协议：

  • 是一套维护缓存一致性的算法
  • 如果所有处理器的缓存控制器都按照缓存一致性协议的规定进行运作，那么缓存一致性就能得到保障

基于失效（invaliditation）的一致性协议

  • 关键思想：
    • 一个处于“modified”状态的缓存行，能够在不通知其他处理器的情况下被修改
    • 处理器只能写入那些处于“modified”状态的缓存行
      • 需要一个方法告诉其他处理器想修改这个缓存行（通常是通过广播消息）
    • 当某个缓存控制器看到了对缓存行的修改请求，它必须使它自己的cahce中对应的缓存行“失效“

MSI缓存控制协议（是一种基于失效的一致性协议）

  • MSI规定了一个缓存行有三种状态：
    • I (Invalid): 失效状态
    • S (shared)：共享状态，可在多个cache中共存，其中的内容是最新的
    • M (Modified)：修改状态，只能在一个cahce中存在
  • 规定了两种处理器操作：
    • PrRd（read）
    • PrWr（write）
  • 三种和缓存一致性有关的总线事务（transaction）
    • BusRd：没有写意图地获取缓存行
    • BusRdx：有写意图地获取缓存行
    • BusWB：将dirty的缓存行写回内存

MSI协议的状态转换图如下所示：
image-20260623001744881

MESI缓存控制协议

  • 为了解决MSI存在的效率问题，额外增加一个state “E”-- exclusive clean
  • image-20260623230424542image-20260623230436142

状态转换图如下：

image-20260623230034356

基于嗅探的一致性协议的可扩展性太差，由此推出了基于directory的控制协议，略过

Lecture 12 Memory Consistency(/Ordering/Model)

这部分内容还参考了CMU15418

Cache coherence V.S. Memory Consistency

  • Cache coherence: 定义了同一内存地址进行读写的行为，依赖多个核对一个内存地址在缓存与内存之间的同步方法

    • 只保证对内存地址X的写操作最终会按照程序顺序正确地回写到内存中
    • 它的目标是保证一个多核多缓存的处理器架构的行为和没有缓存的系统一致

  • Memory Consistency：定义对多个内存地址的读写行为（相对于其他处理器而言）。

    • 处理内存地址X什么时候写回到内存中的问题（相对于其他读写操作而言）

Sequential Consistency

什么是Sequential Consistency？直觉上理解：

  1. 每个处理器对内存的读取和写入都以program order进行
  2. 所有对内存的存取操作都是顺序进行的

举例如下图，内存有A、B两个变量，初始值都为0，处理器0执行 A = 1；Ready = 1，处理器1执行x=ready；y = A。那么所有可能的结果（如果按照sequential Consistency）中除了(1,0)以外所有其他结果都有可能。这是因为如果x=1，说明P0已经执行完成了(b),按照以上”直觉上的理解“，说明P0也已经将(a)执行完成了，那么y只能等于1。反过来说，如果x=1但y=0，说(b)和(a)的执行顺序被对调了，这不符合我们对Sequential Consistency的预期。
image-20260701224822024

具体来讲，Sequential Consistency遵循四种内存序：

  • Wx -> Ry : 对x的写必须发生在对y读之前
  • Rx -> Ry: 对x的读必须发生在对y读之前
  • Rx->Wy :...
  • Wx->Wy: ...

但是为了提升性能，可以不用遵循所有四种内存序，

  1. TSO（TotalStoreOrdering）：不遵循Wx -> Ry的顺序规定

    image-20260702051537207

  2. PSO(Partial Strore Ordering):在TSO的基础上，不遵循Wx->Wy的顺序。这种情况下，上面例子中的x，y值就有可能是(1,0)

    image-20260702051903713

  3. WO(WeakOrdering)/RC(ReleaseConsistency): 所有四种内存序都不遵循：

    image-20260702052120741

但在TSO和WO下如何保证程序的正确性呢？ —— 正确性是指程序在TSO或者WO下的程序运行结果与squential Consistency下的一样

  • 识别出data race的部分，并将他们同步，程序就有了SC一样的结果：

    image-20260701230549586

  • 由此诞生了DRF-SC（Data Race Free -- Sequential Consistency），也就是说程序员只要写出无数据竞争的代码，那么程序运行结果就是和SC内存模型的结果一致，尽管程序运行的硬件使用更加宽松的内存模型

  • 为了支持同步操作，各硬件平台提供支持，比如x86平台有 mm_lfence（wait for all load operations） mm_sfence（wait for all store operations） mm_mfence（wait for all mem operations）等指令

    • 语言级别的同步操作，比如lock、barrier、atomic等操作自动地帮助我们插入硬件平台的同步指令。因此这就能运行程序员，能在一个更高层抽象上保证程序无数据竞争，也就保证了与SC结果的一致性
    • 具体的例子：atomic 变量或 atomic 操作更贴切的叫法其实是 synchronizing atomics，顾名思义它的作用其实有两个（lock、barrier也是类似的）
      1. 防止多线程读写的数据竞争
      2. 插入同步指令，维持某种内存序

内存模型总结：

  1. 定义了几种被打乱的内存序（由编译器或者硬件打乱——更多的是被硬件打乱，编译器的打乱可能导致更松散的内存模型，可参考这里的第二篇文章）
  2. 为什么有宽松的内存模型？--- 因为性能
  3. 为了保证应用程序的正确性，编译器/硬件 与 应用程序达成了某种契约（即DRF-SC），只要应用程序按照DataRaceFree的方式去写，那么整个程序的结果就是SC程序的结果，尽管编译器/硬件使用更加宽松的内存序
  4. 内存模型的细节被隐藏在了 synchronization里了，程序员通常不感知（但让也有语言提供了这种感知以及某种程度的控制，比如c++的std::memory_order）

其他推荐阅读：

research!rsc: Memory Models

C++11中的内存模型下篇 - C++11支持的几种内存模型 | codedump notes

posted @ 2026-06-28 22:46  别杀那头猪  阅读(4)  评论(0)    收藏  举报