Skip to content

全部文章

第 33 頁,共 129 頁

詞表怎麼選、token 怎麼吐,以及 speculative decoding 為什麼快|CMU 11-868 第 8–9 堂

L08 從 BPE 講到講者 Lei Li 參與的 VOLT:詞表大小有成本也有價值,VOLT 用「每加一個 token 能降多少正規化熵」找划算的點,再化成最佳傳輸問題求解;後半講 LLaMA 3 詞表從 32k 擴到 128k、中文字被 byte-level BPE 切成三個 token 的代價。L09 從 greedy、取樣、beam search 一路講到 speculative decoding:小模型先猜 N 個 token,大模型一次前向驗證,因為驗證比生成便宜;最後介紹 EAGLE 改成預測最後一層特徵。

TPU、JAX 與 Pallas:同一個 attention,在 TPU 上從 XLA 融合寫到 Splash Attention|CMU 11-868 第 12–13 堂

Google 的 Srinath Mandalapu 用兩講、兩百多頁,把一個 attention 從 Python 一路追到 TPU 的 VLIW 指令。L12 講 JAX 生態、TPU Ironwood 的記憶體與運算單元,以及 XLA 怎麼把 attention 編成三個融合 kernel。L13 講 XLA 做不到的部分:用 Pallas 自己控制 HBM 與 VMEM 之間的搬運,寫出 FlashAttention,再加上區塊稀疏變成 Splash Attention。思路跟 GPU 相同,差別在於:TPU 上排程主要交給編譯器,要介入就得用 Pallas 把迴圈與區塊大小攬回自己手上。

用系統的眼光讀 Transformer、T5、LLaMA 與 GPT-3|CMU 11-868 第 6–7 堂

11-868 只花兩堂課講模型本身:L06 把 Transformer 拆成 embedding、多頭注意力、FFN、LayerNorm 與殘差,L07 用 T5、LLaMA、GPT-3 三個模型示範現代 LLM 改了哪些地方。對系統工程師來說,重點是記住形狀:GPT-3 175B 是 96 層、d_model 12288、context 2048,訓練 3,000 億 token;LLaMA 65B 是 80 層、d_model 8192、訓練 1.4 兆 token。這些數字決定了後面每一堂加速、平行與服務課要處理的量。

ZeRO 怎麼把資料平行的記憶體切掉——optimizer state、梯度、參數各切一刀|CMU 11-868 第 18 堂

資料平行每張卡都存一份完整的參數、梯度和 optimizer state;用 Adam 混合精度訓練時,每個參數約要 20 bytes,其中 16 bytes 是 optimizer 相關,LLaMA-3 8B 光這些就要 160GB。CMU 11-868 L18 以 ZeRO 論文為主軸,逐格動畫示範三個階段:ZeRO-1 切 optimizer state、ZeRO-2 再切梯度、ZeRO-3 連參數也切。投影片的結論是前兩階段不增加通訊、最多省 8 倍記憶體;第三階段每張卡的用量隨 GPU 數下降,代價是投影片估的約 3 倍通訊。

從一顆晶片到整座資料中心——資料流硬體、kernel 融合、平行策略與記憶體瓶頸|Stanford CS149 第 12 堂

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 都是在解同一個問題。這講沒有公開錄影,本文只依投影片。

Cache coherence:MSI、MESI 與 false sharing|Stanford CS149 第 14 堂

每個核心都有自己的 cache,同一個位址就可能同時有好幾份副本,各核看到的值會不一樣。這不是加鎖能解決的問題,是硬體複製資料造成的。CS149 L14 先定義什麼叫 coherent,再拆解 snooping 的 MSI 協定:要寫就先廣播 BusRdX 讓別人作廢,MESI 多一個 E 狀態省掉「讀完再寫」的第二筆交易,directory 則把廣播改成點對點。對程式設計師最實際的後果是 false sharing:兩個 thread 寫不同變數,只因為落在同一條 cache line,就讓 cache line 在核心間來回彈,投影片的 demo 慢了三倍。

平行計算:課程定位與兩條閱讀路線|Stanford CS149 課程總覽

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 版,本系列只拿它當聽講補充。

資料平行思維,用 map、scan 與 sort 取代鎖|Stanford CS149 第 8 堂

第 8 講要你換一個腦袋:不再想「每個 worker 做什麼」,而是把演算法寫成對序列的操作,例如 map、fold、scan、segmented scan、gather/scatter、sort、groupBy。這些原語都有高效的平行實作,能把不規則的平行變規則、把細粒度同步變粗粒度。代價是要多掃幾遍資料,所以很吃頻寬。

在 GPU 上高效跑 DNN——卷積改寫成矩陣乘、分塊、融合,一路推到 FlashAttention|Stanford CS149 第 9 堂

L9 開場就說:懂了 arithmetic intensity 與 roofline,你就幾乎懂了現代 AI 軟體面效能最佳化的一切。接著示範三件事:全連接層、卷積層、attention 最後都變成矩陣乘(GEMM);GEMM 要靠分塊(blocking)讓資料留在快取裡;層與層之間要融合(fusion),別把中間結果寫回 DRAM 再讀回來。softmax 可以分塊計算,這就是融合版 attention(FlashAttention 的核心想法)不必存下 N×N 矩陣的原因。

讓效能最佳化不再只靠專家——Halide 的演算法/排程分離、自動排程器與 LLM agent|Stanford CS149 第 13 堂

L13 問的是:寫快程式的專家太少,怎麼辦?投影片給三個答案。一是提高抽象層級:Halide 把「算什麼」(演算法)和「怎麼算」(排程)拆成兩種語言,同一段模糊濾波可以只改一行排程就換成分塊、向量化、多核版本。二是智慧搜尋:因為排程空間定義清楚,可以用搜尋加學出來的成本模型自動產生排程。三是新興的 LLM agent:讓模型寫 CUDA、執行、看 profiler、反思、再改,並用範例資料庫或 prompt 最佳化讓它自我改進。投影片最後把問題留給你:真正的價值在 DSL 設計,還是 LLM agent?

細粒度鎖與 lock-free:從 test-and-set 到 ABA 問題|Stanford CS149 第 16 堂

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 的系統。

GPU 架構與 CUDA 程式設計|Stanford CS149 第 7 堂

CUDA 的 grid、thread block、CUDA thread 是一套程式抽象;GPU 用 SM、warp 與硬體 block 排程器把它實作出來。這一講的核心是分清兩件事:thread block 之間系統可以任意排序,同一個 block 裡的 thread 則保證同時存在,所以 block 能用 shared memory 和 __syncthreads() 合作,也所以一個 SM 能塞幾個 block 由暫存器與 shared memory 的量決定。

通用處理器為什麼浪費能量——硬體專用化、Tensor Core、TPU 脈動陣列與資料流架構|Stanford CS149 第 10 堂

L10 從一條式子出發:功耗固定時,效能只能靠「每個運算花多少焦耳」來換,而通用處理器把大部分能量花在取指令、解碼、搬資料,不是在算。投影片的經驗法則是 GPU 比 CPU 好約 10 倍 perf/watt,固定功能 ASIC 可達 100–1000 倍。接著用同一套標準(分塊 tensor、非同步計算與記憶體、運算單元直接互傳)檢視 H100 的 Tensor Core 與 TMA、Google TPU 的脈動陣列,以及可重組資料流架構。

處理器很快,資料送不過來——延遲 vs 頻寬,以及 ISPC 教你分清「抽象」與「實作」|Stanford CS149 第 3 堂

L3 前半用高速公路與洗衣服的比喻分開延遲與頻寬,再算給你看:向量逐元素相乘在 V100 上效率不到 1%,因為記憶體送資料的速度追不上 ALU。後半的主題是「抽象 vs 實作」:ISPC 讓你用 SPMD 的方式思考(一群 program instance 各做一份),編譯器卻用 SIMD 指令實作。把這兩層混在一起,是這門課最常見的困惑來源。

Locality、通訊與 arithmetic intensity:為什麼少搬資料比多開核更重要|Stanford CS149 第 6 堂

CS149 L6 要你把「通訊」看得很廣:處理器和 cache、和記憶體、和另一台機器之間的資料搬移都算。現代平行處理器的算力遠高於頻寬,所以 arithmetic intensity(每搬一單位資料做多少計算)決定了你能不能把硬體餵飽。提高它的手段有三類:改分配方式減少必要通訊、用 blocking 與 loop fusion 減少 cache 造成的額外通訊、用分散與錯開存取降低 contention。

多核、SIMD、硬體多執行緒,三種平行各解決什麼問題|Stanford CS149 第 2 堂

CS149 Fall 2025 第二講用一個計算 sin(x) 的迴圈,依序加上三個想法:把電晶體拿去做更多核心(multi-core)、讓一道指令同時驅動多個 ALU(SIMD)、在同一個核心上交錯執行多個執行緒來隱藏記憶體延遲(hardware multithreading)。前兩個增加運算能力,第三個讓運算單元在等記憶體時不閒著。結論是三個條件:平行工作要夠多、同一組工作要跑同樣的指令、平行工作要比 ALU 更多才能隱藏延遲。

在四核 CPU 上量 speedup,並解釋它為什麼不是線性|Stanford CS149 程式作業 1/書面作業 1

PA1 程式寫得不多,重點是分析:六個程式分別練 threads 的工作分配、SIMD 遮罩、ISPC gang 與 task、輸入資料如何影響 SIMD 效率、頻寬受限的 saxpy,以及用計時找 K-Means 的熱點。Written 1 則用紙筆題練峰值吞吐量、指令相依、管線、多執行緒藏延遲與 SIMD divergence。官方評分以 Stanford myth 機器為準,校外可以在自己的機器跑,但數字不能直接對照。

從零打造 task execution library——thread pool、sleep 到有依賴的 task graph|Stanford CS149 程式作業 2

CS149 PA2 要你用 C++ 寫一個多核 CPU 上的任務執行庫,而且要寫四次:每次 run() 都開 thread、改成 spinning 的 thread pool、改成會睡覺的 thread pool,最後擴充成非同步、有依賴的 task graph。每一步都得是完全正確的系統,並跟官方參考實作比速度。官方評分機器是 AWS c7g.4xlarge;校外可以在自己的多核機器上跑,但數字不能直接和官方比。

與 Written 2:CUDA 圓形渲染器,順序正確與速度要一起拿|Stanford CS149 程式作業 3/書面作業 2

PA3 有三部分:把 SAXPY 改寫成 CUDA 並分開計時、用 exclusive scan 實作 find_repeats、再寫一個又對又快的 CUDA 圓形渲染器(85 分)。渲染器的難點是半透明圓的混色不可交換,每個像素都得照輸入順序更新,而起始程式碼一個圓一個 thread 的做法兩樣都沒守住。Written 2 則是五題計分題(fusion、SIMD 利用率、用 barrier 取代鎖、用資料平行原語處理圖、粒子模擬的鎖)加 14 題練習。校外要自備 NVIDIA GPU,本文不提供解答。

在 Trainium2 上自己搬資料——NKI、SBUF/PSUM 與 fused conv+maxpool|Stanford CS149 程式作業 4/書面作業 3

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。