程式開發環境與未來展望 (Programming Environment and Future Outlook)

重點總覽 (Overview)

統一裝置記憶體位址空間 (Unified Device Memory Space)

• 早期 CUDA 裝置中,shared / local / global memory 各自擁有獨立的位址空間;開發者只能對 global memory 使用指標,其他不行。
• 從 Fermi (2009) 起,這些記憶體被併入一個統一位址空間 (unified address space):
  • 單一組 load/store 指令 + 單一種指標即可存取 global / local / shared 任一空間(不再需要 ld.global / ld.shared / ld.local 等不同指令)。
  • 抽象掉「運算元位於哪種記憶體」這件事,開發者只在 allocation 時關心放在哪裡。

[Fermi 之前] 各自獨立的位址空間 — 一個 device function 需為每種記憶體寫一份實作
  ptr_g ──ld.global──> [ global memory ]
  ptr_s ──ld.shared──> [ shared memory ]
  ptr_l ──ld.local ──> [ local  memory ]

[Fermi 之後] 統一位址空間 — 單一 ld/st 指令、單一 device function 即可
                    ┌───────────────────────────┐
        ptr ──ld──> │  single virtual addr space│
                    │   ├── global  (慢)        │
                    │   ├── local             │
                    │   └── shared  (快)        │
                    └───────────────────────────┘

• 可組合性 (composability):同一個 device function 可接受指向任一種 GPU 記憶體的指標參數。
  • 沒有統一位址空間時,一個 device function 得為每種記憶體型態各寫一份實作。

// 統一位址空間下:同一份 device function 不論 p 指向
// global / shared / local memory 都能運作 (composable)
__device__ float reduce(float* p, int n) {
    float s = 0.0f;
    for (int i = 0; i < n; ++i) s += p[i];   // 同一組 ld 指令
    return s;
}

• 額外效益:
  • 大幅降低打造 production 級 CUDA 函式庫的成本。
  • 啟用完整 C / C++ 指標支援(當時的重大進展)。
  • kernel 中已支援 new、delete、constructor、destructor;未來編譯器將強化 C++ templates 與 virtual function call 支援。

高階程式設計介面 (High-Level Programming Interfaces)

• 不斷演進的程式介面持續提升異質平行開發者的生產力,讓更多應用能以最小開發成本取得合理效能。

// OpenACC:在循序迴圈上加 directive,編譯器產生 CUDA kernel
#pragma acc kernels
for (int i = 0; i < n; i++)
    y[i] = a * x[i] + y[i];          // SAXPY → 自動 GPU 化

// Thrust:描述運算,函式庫自動生成/設定 kernel
#include <thrust/device_vector.h>
#include <thrust/sort.h>
thrust::device_vector<int> d(h.begin(), h.end());
thrust::sort(d.begin(), d.end());    // 無須手寫 sort kernel

關鍵路徑剖析 (Profiling with Critical Path Analysis)

• 在 CPU/GPU 都有大量運算的異質應用中,找對優化位置很難:理想是以最少努力換最大 speedup。
• PC sampling (CUDA 7.5):提供指令級 (instruction-level) profiling,讓使用者精準定位最耗時的程式行。
• 但有陷阱:執行時間最長的 kernel 不一定是最該優化的對象。

Fig 22.2:為何要看 critical path

時間 ───────────────────────────────────────────────►
CPU :  |======= A =======|···idle···|== 使用 Y 的結果 ==|
GPU :  |==== Kernel X ====|                              (與 A 完全重疊)
GPU :       |======== Kernel Y ========|                 (位於 critical path)
                                       ▲
                              CPU 在此空轉等待 Y 完成

• Kernel X 較長,但執行時間「完全被 CPU 活動 A 重疊」
  → 只加速 X 而不同時加速 A,應用幾乎不會變快 (speedup ≈ 1)
• Kernel Y 較短,卻在 critical path 上;CPU 空等 Y
  → 加速 Y 可減少 CPU 等待時間 ⇒ 才是最佳優化目標

• Visual Profiler (CUDA 8, 2016):提供 GPU kernel 與 CPU CUDA API call 之間的 critical path analysis。
  • 不在關鍵路徑上的 kernel / memory copy / API call 會被灰階淡化 (grayed out);只有關鍵路徑上的活動以顏色高亮,讓使用者一眼鎖定優化標的。

未來展望 (Future Outlook)

• CUDA 的演進持續強化開發者生產力與現代軟體工程實踐的支援。
• 隨新能力與程式語言支援到位,「以最小開發成本取得合理效能」的應用範圍將大幅擴展。
• 開發者已實際體驗到相較舊版 CUDA 的開發、移植、維護成本下降。
• 以 Thrust 等高階工具自動生成 CUDA 程式碼的既有應用,常能在新版立即獲得效能提升(無須改碼)。

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 22-Advanced-Practices-And-Future-Evolution/01-Host-Device-Interaction-Memory-Model
• 22-Advanced-Practices-And-Future-Evolution/02-Kernel-Execution-Control
• 22-Advanced-Practices-And-Future-Evolution/03-Memory-Bandwidth-and-Compute-Throughput
• 05-Memory-Architecture-And-Data-Locality/01-Memory-Access-Efficiency-and-CUDA-Memory-Types
• 19-Parallel-Programming-And-Computational-Thinking/01-Goals-of-Parallel-Computing-and-Amdahls-Law
• 23-Conclusion-And-Outlook/02-Future-Outlook