Skip to content
所有標籤

#performance

30 篇文章

CMU 11-868 L02–L04 GPU 程式模型與加速:thread、block、記憶體階層與 tiling

CMU 11-868 的 GPU 三講回答一個問題:為什麼寫對的 CUDA matmul 在 A100 上只用到 2.48% 的 FP32 算力。L02 講 SM、warp 與 grid/block/thread,L03 講 cudaMalloc、cudaMemcpy 與 kernel 索引,L04 用 tiling、coalesced access 與避開 bank conflict,把資料從約 500 個 cycle 外的 global memory 搬近一點。

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

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

Stanford CS149 導讀:Fall 2025 平行計算課程總覽

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 L9:在 GPU 上高效跑 DNN——卷積改寫成矩陣乘、分塊、融合,一路推到 FlashAttention

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

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

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

CS149 L16 細粒度鎖與 lock-free:從 test-and-set 到 ABA 問題

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

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

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

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

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

CS149 L6 Locality、通訊與 arithmetic intensity:為什麼少搬資料比多開核更重要

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

CS149 L2:多核、SIMD、硬體多執行緒,三種平行各解決什麼問題

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

CS149 PA1 + Written 1:在四核 CPU 上量 speedup,並解釋它為什麼不是線性

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

CS149 PA2:從零打造 task execution library——thread pool、sleep 到有依賴的 task graph

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

CS149 PA4 + Written 3:在 Trainium2 上自己搬資料——NKI、SBUF/PSUM 與 fused conv+maxpool

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。

CS149 PA5 在 H100 上寫最快的 kernel:五道 AI kernel 題,評分看工作日誌不看排名

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 跑。

CS149 L4:拿到一個程式,怎麼把它平行化?——分解、分配、協調,以及 Amdahl 定律

L4 給了一套平行化的思考流程:先分解(decomposition)找出彼此獨立的工作,再分配(assignment)給工作者,再協調(orchestration)通訊與同步,最後對應到硬體。Amdahl 定律提醒你循序部分決定 speedup 上限。貫穿全講的例子是 2D 網格解算器:原本的相依關係很難平行,改用紅黑著色換一種更新順序後,就能用資料平行或共享位址空間兩種模型寫出來。

CS149 L11:怎麼替專用硬體寫程式——ThunderKittens 管住 H100 的非同步,資料流架構用 metapipeline 換掉它

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。

CS149 L1:單核為什麼不再變快,以及「快」為什麼不等於「有效率」

CS149 Fall 2025 第一講先定義 speedup,再用三個課堂示範說明通訊與負載不均會吃掉加速比。接著解釋單核效能為什麼停滯:superscalar 能挖的指令層級平行大約在每時脈發四道指令就用完,時脈又被功耗牆擋住,所以效能只能靠多核與專用硬體。最後一段把焦點轉到效率:取一次 DRAM 的延遲約是 L1 cache 的 60 倍,搬 64 bits 的能耗是一次整數運算的上千倍,高效率幾乎都歸結到高效率地存取資料。

CS149 L5 工作分配與排程:從 work queue 到 Cilk 的 work stealing

負載平衡的難處在於它和排程成本互相拉扯:任務切得越細越好平衡,但每次領任務都要付同步成本。CS149 L5 先把選項排成一條從 static 到 dynamic 的連續光譜,再拆解 Cilk 的執行期:每個 worker 一條 deque,spawn 時先跑 child、把 continuation 留給別人偷,閒置的 thread 從別人 deque 的頂端偷走最大塊的工作。

aiguideMIT 6.5940 導讀

MIT 6.5940 Fall 2026 Lab 1 補充:用 Roofline、Profiler 與 FlashAttention 看懂 GPU 瓶頸

這篇是 Fall 2026 的材料,不屬於本系列主幹的 Fall 2024。Fall 2026 把 Lab 1 從剪枝換成「Efficient AI Fundamentals」(lab1_gpu_basics.zip):Part 1 手刻三層迴圈的 GEMM、算 MAC/FLOPs/I/O;Part 2 畫 GEMM 與 GEMV 的 roofline;Part 3 拿 gemma-3-270m-it 的 decoder layer 算 attention 與 MLP,比較 prefill 與 decode;Part 4 用 PyTorch Profiler 看 kernel、自己寫 GeLU 體會 kernel fusion、再試 torch.compile 與 CUDA Graph;Part 5 比較 SDPA 與 FlashAttention。主題 80 分,另有 20 分 bonus,Part 5 因為 Colab T4 跑不動整段改成 bonus。

aiguideMIT 6.5940 導讀

MIT 6.5940 L11 TinyEngine 與平行運算:從 loop tiling 到 in-place depthwise

演算法把模型縮小之後,系統層還能再榨多少?L11 用同一個矩陣乘法示範:loop reordering 快 12 倍、tiling 快 19 倍(Intel Xeon 4114),CUDA 版在 2080Ti 上端到端快 94 倍。後半講 TinyEngine 用的推論技巧:im2col、in-place depthwise 把峰值記憶體從 2×C×H×W 降到 (1+C)×H×W、pointwise 用 NHWC、depthwise 用 NCHW,以及少 2.25 倍乘法的 Winograd。

aidebug

每修一次都以為好了:Vision PDF 解析的三階段優化

用 Vision API 解析 150 頁 PDF 要 29 分鐘(一頁一請求)。分批送頁降到 1.5 倍,加上 5 路併發壓到 24 秒。上線一週後客戶一次傳 23 份 PDF,70 個並行請求打爆 Bedrock 限流——90 頁被跳過、5 份失敗、8 份卡住。最終用 Redis slot 做全叢集上限 + backoff 重試修復。

Cloudflare Cache Rules:什麼該快取,什麼要保持動態

Cloudflare Cache Rules 是 zone 層的快取政策:用 request expression 決定哪些內容 eligible for cache、Edge TTL/Browser TTL 怎麼算、cache key 包哪些維度,以及遇到 stale、ETag、purge 時怎麼處理。它適合管理 CDN 快取規則;Worker Cache API 則適合程式化存取。

Cloudflare Images 怎麼用:圖片變體、格式轉換與交付管線

Cloudflare Images 有兩條路:把 R2/S3/origin 圖片拿來做 edge transformations,或直接把圖片存進 Images 再用 variant delivery。前者按 unique transformations 看成本,後者還要看 stored 和 delivered images。

Cloudflare Smart Shield 怎麼用:讓 Origin 少扛一點流量

Smart Shield 是 Cloudflare 的 origin protection bundle:用 Smart Tiered Cache、connection reuse、Argo Smart Routing、Regional Tiered Cache、Cache Reserve、Health Checks 和 Dedicated CDN Egress IPs,減少打到 origin 的 request 與 connection。

CS336 Lecture 5:GPU 快不是因為每個 thread 快,而是資料少搬幾次

第五講從 SM、warp 與記憶體階層解釋 GPU,再用低精度、fusion、recomputation、coalescing 與 tiling 統一常見優化;FlashAttention 正是這些原則在 attention 上的組合。

CS336 Lecture 6:寫 Triton kernel 前,先學會 benchmark 與 profile

第六講把 GPU 原理落到 kernel:benchmark 看不同尺寸如何縮放,profiler 看實際呼叫與時間,再以 Triton 實作 GeLU、softmax、reduction 與 tiled matmul;快的前提是先量對。

CS336 Lecture 2:先算 FLOPs 與記憶體,再談模型跑不跑得動

第二講把模型訓練還原成 tensor、FLOPs、bytes 與時間:用 einops 管維度,以 arithmetic intensity 和 roofline 判斷瓶頸,再用 gradient accumulation 與 activation checkpointing 交換運算和記憶體。

Stanford CS107 Lecture 25:Caching、Memory Hierarchy 與 Locality

CS107 第 25 講用精簡投影片建立 cache 的核心模型:記憶體存取成本不均,較小且較快的層級保存可能再次使用的資料,而 temporal 與 spatial locality 決定程式能否受益。

aiguideRAG 技法大全

RAG 成本優化:把每次查詢的花費壓到最低

RAG 系統的成本來自 LLM token、Embedding API、向量搜尋。每個環節都有可以壓成本的地方,但要確認優化沒有犧牲太多品質。

aiguideRAG 技法大全

Semantic Caching:語義相近的問題只跑一次 RAG

快取不只能比對完全一樣的查詢,語義相近的問題也能命中快取,省下整個 RAG pipeline 的執行。