系列CS149 是 Kayvon Fatahalian 與 Kunle Olukotun 在 Stanford 教的平行計算課,從多核 CPU、SIMD 一路講到 GPU、AI 加速器、資料中心,最後回到 cache coherence 與 lock-free。Fall 2025 的 18 份投影片、5 個程式作業的 starter code 與 README、4 份書面作業 PDF 都能匿名取得,本系列評為 A3(足以自學)。缺口有四個:當期錄影只在 Canvas;PA1 評分用 Stanford myth 機器;PA4 要自費租 AWS Trainium2,而且課程 AMI 是私有的;PA5 的 H100 排隊系統與排行榜要 SUNet ID。公開錄影是 2023 版,本系列只拿它當聽講補充。
CS149 Fall 2025 第一講先定義 speedup,再用三個課堂示範說明通訊與負載不均會吃掉加速比。接著解釋單核效能為什麼停滯:superscalar 能挖的指令層級平行大約在每時脈發四道指令就用完,時脈又被功耗牆擋住,所以效能只能靠多核與專用硬體。最後一段把焦點轉到效率:取一次 DRAM 的延遲約是 L1 cache 的 60 倍,搬 64 bits 的能耗是一次整數運算的上千倍,高效率幾乎都歸結到高效率地存取資料。
CS149 Fall 2025 第二講用一個計算 sin(x) 的迴圈,依序加上三個想法:把電晶體拿去做更多核心(multi-core)、讓一道指令同時驅動多個 ALU(SIMD)、在同一個核心上交錯執行多個執行緒來隱藏記憶體延遲(hardware multithreading)。前兩個增加運算能力,第三個讓運算單元在等記憶體時不閒著。結論是三個條件:平行工作要夠多、同一組工作要跑同樣的指令、平行工作要比 ALU 更多才能隱藏延遲。
L3 前半用高速公路與洗衣服的比喻分開延遲與頻寬,再算給你看:向量逐元素相乘在 V100 上效率不到 1%,因為記憶體送資料的速度追不上 ALU。後半的主題是「抽象 vs 實作」:ISPC 讓你用 SPMD 的方式思考(一群 program instance 各做一份),編譯器卻用 SIMD 指令實作。把這兩層混在一起,是這門課最常見的困惑來源。
PA1 程式寫得不多,重點是分析:六個程式分別練 threads 的工作分配、SIMD 遮罩、ISPC gang 與 task、輸入資料如何影響 SIMD 效率、頻寬受限的 saxpy,以及用計時找 K-Means 的熱點。Written 1 則用紙筆題練峰值吞吐量、指令相依、管線、多執行緒藏延遲與 SIMD divergence。官方評分以 Stanford myth 機器為準,校外可以在自己的機器跑,但數字不能直接對照。
L4 給了一套平行化的思考流程:先分解(decomposition)找出彼此獨立的工作,再分配(assignment)給工作者,再協調(orchestration)通訊與同步,最後對應到硬體。Amdahl 定律提醒你循序部分決定 speedup 上限。貫穿全講的例子是 2D 網格解算器:原本的相依關係很難平行,改用紅黑著色換一種更新順序後,就能用資料平行或共享位址空間兩種模型寫出來。
負載平衡的難處在於它和排程成本互相拉扯:任務切得越細越好平衡,但每次領任務都要付同步成本。CS149 L5 先把選項排成一條從 static 到 dynamic 的連續光譜,再拆解 Cilk 的執行期:每個 worker 一條 deque,spawn 時先跑 child、把 continuation 留給別人偷,閒置的 thread 從別人 deque 的頂端偷走最大塊的工作。
CS149 L6 要你把「通訊」看得很廣:處理器和 cache、和記憶體、和另一台機器之間的資料搬移都算。現代平行處理器的算力遠高於頻寬,所以 arithmetic intensity(每搬一單位資料做多少計算)決定了你能不能把硬體餵飽。提高它的手段有三類:改分配方式減少必要通訊、用 blocking 與 loop fusion 減少 cache 造成的額外通訊、用分散與錯開存取降低 contention。
CS149 PA2 要你用 C++ 寫一個多核 CPU 上的任務執行庫,而且要寫四次:每次 run() 都開 thread、改成 spinning 的 thread pool、改成會睡覺的 thread pool,最後擴充成非同步、有依賴的 task graph。每一步都得是完全正確的系統,並跟官方參考實作比速度。官方評分機器是 AWS c7g.4xlarge;校外可以在自己的多核機器上跑,但數字不能直接和官方比。
CUDA 的 grid、thread block、CUDA thread 是一套程式抽象;GPU 用 SM、warp 與硬體 block 排程器把它實作出來。這一講的核心是分清兩件事:thread block 之間系統可以任意排序,同一個 block 裡的 thread 則保證同時存在,所以 block 能用 shared memory 和 __syncthreads() 合作,也所以一個 SM 能塞幾個 block 由暫存器與 shared memory 的量決定。
第 8 講要你換一個腦袋:不再想「每個 worker 做什麼」,而是把演算法寫成對序列的操作,例如 map、fold、scan、segmented scan、gather/scatter、sort、groupBy。這些原語都有高效的平行實作,能把不規則的平行變規則、把細粒度同步變粗粒度。代價是要多掃幾遍資料,所以很吃頻寬。
PA3 有三部分:把 SAXPY 改寫成 CUDA 並分開計時、用 exclusive scan 實作 find_repeats、再寫一個又對又快的 CUDA 圓形渲染器(85 分)。渲染器的難點是半透明圓的混色不可交換,每個像素都得照輸入順序更新,而起始程式碼一個圓一個 thread 的做法兩樣都沒守住。Written 2 則是五題計分題(fusion、SIMD 利用率、用 barrier 取代鎖、用資料平行原語處理圖、粒子模擬的鎖)加 14 題練習。校外要自備 NVIDIA GPU,本文不提供解答。
L9 開場就說:懂了 arithmetic intensity 與 roofline,你就幾乎懂了現代 AI 軟體面效能最佳化的一切。接著示範三件事:全連接層、卷積層、attention 最後都變成矩陣乘(GEMM);GEMM 要靠分塊(blocking)讓資料留在快取裡;層與層之間要融合(fusion),別把中間結果寫回 DRAM 再讀回來。softmax 可以分塊計算,這就是融合版 attention(FlashAttention 的核心想法)不必存下 N×N 矩陣的原因。
L10 從一條式子出發:功耗固定時,效能只能靠「每個運算花多少焦耳」來換,而通用處理器把大部分能量花在取指令、解碼、搬資料,不是在算。投影片的經驗法則是 GPU 比 CPU 好約 10 倍 perf/watt,固定功能 ASIC 可達 100–1000 倍。接著用同一套標準(分塊 tensor、非同步計算與記憶體、運算單元直接互傳)檢視 H100 的 Tensor Core 與 TMA、Google TPU 的脈動陣列,以及可重組資料流架構。
L11 問的是:硬體為 AI 專用化之後,程式設計師要付出什麼?投影片以 H100 為例:要吃滿 Tensor Core,就得用 16×16 tile、讓 TMA 非同步搬資料、讓 producer 與 consumer warp 管線化,寫起來很複雜,所以有了 ThunderKittens 這類 DSL。另一條路是資料流架構(SambaNova SN40L):用 map/reduce/zip 等平行模式描述計算,編譯器做分塊、metapipelining 與佈局,投影片說它能把 Llama 3.1 8B 的整個 decoder 融成一個 kernel。
PA4 把你丟到 AWS Trainium2 的一顆 NeuronCore 上。這裡沒有 cache 幫你決定什麼資料留在晶片上:SBUF(28 MiB)與 PSUM(2 MiB)都要你用 dma_copy 明確搬進搬出,partition 維度最多 128。Part 1 用 vector add 和 transpose 教你這些限制與 DMA 成本,Part 2 要你把 convolution 拆成一串 matmul、再和 max pool 融合成一個不落地到 HBM 的 kernel。Written 3 用 line buffer、兩次 box blur、softmax 硬體與 metapipelining 練同一件事:把中間結果留在晶片上。環境需要課程 private AMI 與付費 capacity block,校外實質只到 A2。
L12 的主軸是「搬資料」。前段用 SambaNova SN40L 說明資料流架構與 metapipelining:同樣跑 Llama 3.1 8B,投影片說 RDU 每個 token 約 3 次 kernel 呼叫,GPU 約 800 次,差別來自能不能把整個 decoder 融合進一個 kernel。中段把尺度拉到資料中心:TP、PP、EP、DP 各自需要哪種集體通訊,以及計算與通訊重疊為什麼決定擴展效率。後段回到能源與 DRAM:搬一個 byte 比算一次貴得多,記憶體控制器、burst mode、HBM 都是在解同一個問題。這講沒有公開錄影,本文只依投影片。
L13 問的是:寫快程式的專家太少,怎麼辦?投影片給三個答案。一是提高抽象層級:Halide 把「算什麼」(演算法)和「怎麼算」(排程)拆成兩種語言,同一段模糊濾波可以只改一行排程就換成分塊、向量化、多核版本。二是智慧搜尋:因為排程空間定義清楚,可以用搜尋加學出來的成本模型自動產生排程。三是新興的 LLM agent:讓模型寫 CUDA、執行、看 profiler、反思、再改,並用範例資料庫或 prompt 最佳化讓它自我改進。投影片最後把問題留給你:真正的價值在 DSL 設計,還是 LLM agent?
CS149 Fall 2025 的最後一份程式作業是開放式的:從 Histogram、1D occupancy decoder、FlashAttention、3D 熱方程 RK4、SwiGLU 五題挑一題以上,在 H100 上把 PyTorch baseline 調快,可以用 CUDA、Triton 或 TileLang,也允許用 LLM。分數不看速度門檻,看你交的工作日誌能不能說清楚每一步量了什麼、推出什麼假設、為什麼停手。H100 job queue 與排行榜要 SUNet ID,校外只能在自己的 NVIDIA GPU 上用 eval.py 跑。
每個核心都有自己的 cache,同一個位址就可能同時有好幾份副本,各核看到的值會不一樣。這不是加鎖能解決的問題,是硬體複製資料造成的。CS149 L14 先定義什麼叫 coherent,再拆解 snooping 的 MSI 協定:要寫就先廣播 BusRdX 讓別人作廢,MESI 多一個 E 狀態省掉「讀完再寫」的第二筆交易,directory 則把廣播改成點對點。對程式設計師最實際的後果是 false sharing:兩個 thread 寫不同變數,只因為落在同一條 cache line,就讓 cache line 在核心間來回彈,投影片的 demo 慢了三倍。
Coherence 只管單一位址;memory consistency 管的是不同位址之間的讀寫,在別的 thread 眼中以什麼順序生效。CS149 L15 用兩個 thread、兩個變數的例子說明:sequential consistency 下 r1 = r2 = 0 不可能,但每顆現代處理器都有的 write buffer 會讓讀跑到寫前面,於是它變成可能。TSO、PSO、weak ordering 依序放寬更多順序換取效能,fence 與同步原語再把需要的順序補回來。對應用程式設計師的結論很短:寫沒有 data race 的程式,用同步函式庫,C11、C++11、Java 5 就保證你看到 sequential consistency。
CS149 L16 分三段:先用 cache coherence 的眼光看鎖怎麼實作(test-and-set、test-and-test-and-set、ticket lock、CAS、LL/SC),再用一條排序 linked list 示範從一把大鎖改成 hand-over-hand 細粒度鎖,最後介紹 lock-free:單一生產者單一消費者佇列、用 CAS 寫的 stack、ABA 問題與 hazard pointer。投影片的結論很務實:在只有你的程式佔用機器的情境下,寫得好的鎖版本常常一樣快,而且好寫得多;lock-free 的價值在於執行緒可能被搶佔、page fault 的系統。
粗鎖好寫但慢,細鎖快但容易寫錯。Transactional memory 讓程式設計師只宣告 atomic { },由系統負責原子性與隔離。CS149 L17 講動機(failure atomicity、composability)與設計空間:資料版本管理分 eager(undo log)與 lazy(write buffer),衝突偵測分 pessimistic 與 optimistic。L18 拆 STM 的執行期資料結構與 McRT 演算法,再講 HTM 怎麼用 cache 的 R/W 位元加上 coherence 協定偵測衝突,最後是 Intel Haswell 的 RTM。Written 4 則把 MSI、LL/SC、鎖與記憶體順序、雙向 linked list 的細粒度鎖串成四題。