Lazy Loading 與 Error Log Management

重點總覽

Lazy Loading 介紹

Lazy loading 透過「等到 CUDA module 真正被需要時才載入」來縮短程式初始化時間。對於只使用所含 kernel 中一小部分的程式特別有效，這在使用 library 時很常見。

• 設計目標：只要遵循 CUDA programming model，lazy loading 對使用者應是「無感」(invisible) 的。
• CUDA 12.3 起所有平台預設啟用，可用 CUDA_MODULE_LOADING 環境變數控制。

eager loading（傳統）            lazy loading（預設）
program start                   program start
   │ 載入全部 module                │ 只初始化，不載入 kernel
   ▼                               ▼
[m0][m1][m2]...[mN] 全載入         launch kernel kX  ──► 此時才載入 mX
   │                               ▼
   ▼                            [mX] 載入
執行                              執行（其餘 module 不佔資源）

Lazy Loading 的需求（Requirements）

Lazy loading 是 CUDA runtime 與 driver 的「聯合功能」，只有同時滿足 runtime 與 driver 版本需求時才可用。

• Runtime 版本：CUDA runtime 11.7 起提供。由於 runtime 通常靜態連結進程式/library，因此只有用 11.7+ toolkit 編譯的程式與 library 能受益；用舊版 runtime 編譯的 library 會 eager 載入所有 module。
• Driver 版本：需 driver 515 或更新版。即使用 11.7+ toolkit，driver 低於 515 也無法使用 lazy loading。
• Compiler 需求：不需要任何 compiler 支援。pre-11.7 編譯的 SASS 與 PTX 都能在 lazy loading 下載入並享有完整效益，但仍需 11.7+ runtime。
• Kernel 需求：lazy loading 不影響含 managed variable 的 module，這類 module 仍會 eager 載入。

Lazy Loading 的使用（Usage）

啟用與停用

• 啟用：CUDA_MODULE_LOADING=LAZY
• 停用：CUDA_MODULE_LOADING=EAGER
• CUDA 12.3 起所有平台預設即為 LAZY。

在執行期檢查是否啟用

用 driver API cuModuleGetLoadingMode 查詢，注意呼叫前 CUDA 必須先初始化。

#include <cuda.h>
#include <assert.h>
#include <iostream>
int main() {
    CUmoduleLoadingMode mode;
    assert(CUDA_SUCCESS == cuInit(0));                       // 必先初始化
    assert(CUDA_SUCCESS == cuModuleGetLoadingMode(&mode));
    std::cout << "CUDA Module Loading Mode is "
              << ((mode == CU_MODULE_LAZY_LOADING) ? "lazy" : "eager")
              << std::endl;
    return 0;
}

重點：先 cuInit(0)，再用 cuModuleGetLoadingMode 取得 mode，比對 CU_MODULE_LAZY_LOADING 判斷 lazy 或 eager。

強制單一 module/kernel 立即（eager）載入

載入 kernel 與變數是自動發生的，不需顯式載入。若想在「不執行」的情況下顯式載入某個 kernel：

• cuModuleGetFunction()：會使 module 被載入到 device memory。
• cudaFuncGetAttributes()：會使 kernel 被載入到 device memory。

Lazy Loading 的潛在陷阱（Potential Hazards）

Lazy loading 設計上不需修改應用程式即可使用，但當程式未完全遵循 CUDA programming model 時會有以下注意事項。

對 concurrent kernel execution 的影響

• 有些程式錯誤地假設 concurrent kernel execution 一定發生。若需要跨 kernel 同步、但 kernel 執行被序列化，就可能 deadlock。
• 緩解方式：
  • 在 launch 前先 preload 所有預期要並行的 kernel；或
  • 用 CUDA_MODULE_LOADING=EAGER 強制 eager 載入資料，但不需逐一強制每個 function 載入。

大型記憶體配置（Large Memory Allocations）

• lazy loading 把 module 的記憶體配置從初始化延後到接近執行時。若程式在啟動時就配置整個 VRAM，CUDA 可能在執行期無法為 module 配到記憶體。
• 可能解法：
  • 改用 cudaMallocAsync()，而非開機即吃滿整片 VRAM 的 allocator。
  • 預留一些緩衝空間，補償 kernel 延後載入的需求。
  • 在初始化 allocator 前，先 preload 程式會用到的所有 kernel。

對效能量測（Performance Measurements）的影響

• lazy loading 可能把 CUDA module 初始化移進「被量測的執行視窗」，使量測結果失真。
• 避免方式：
  • 量測前至少做一次 warmup 迭代。
  • 在 launch 前 preload 要被 benchmark 的 kernel。

Error Log Management

Error Log Management 機制讓 CUDA API 錯誤能以「白話英文」(plain-English) 格式回報，描述問題成因。

Background（背景）

• 傳統上，CUDA API 呼叫失敗的唯一指示就是回傳一個非零碼。
• CUDA Toolkit 12.9 起，CUDA Runtime 為錯誤情況定義了 100 多種回傳碼，但許多碼很「通用」(generic)，對除錯成因毫無幫助。

Activation（啟用）

• 設定 CUDA_LOG_FILE 環境變數，可接受值為 stdout、stderr，或系統上一個合法的檔案寫入路徑。
• 即使執行前未設 CUDA_LOG_FILE，仍可透過 API dump log buffer。
• 無錯誤的執行可能完全不會印出任何 log。

Output（輸出格式）

Log 輸出格式如下：

[Time][TID][Source][Severity][API Entry Point] Message

例如，當開發者嘗試把 Error Log 的 logs dump 到一個未配置的 buffer 時，會產生如下實際錯誤訊息：

[22:21:32.099][25642][CUDA][E][cuLogsDumpToMemory] buffer cannot be NULL

對比之下，過去開發者只會得到回傳碼 CUDA_ERROR_INVALID_VALUE，若呼叫 cuGetErrorString 也頂多得到 “invalid argument”。

API Description（API 說明）

CUDA Driver 提供兩類 API 與 Error Log Management 互動：callback 註冊與log 輸出管理。

callback 簽章與註冊/反註冊：

// callback 簽章：每當產生一筆 error log 就被呼叫
void callbackFunc(void *data, CUlogLevel logLevel, char *message, size_t length);

// 註冊；userData 原封不動傳給 callback；callback_out 須由呼叫者保存供反註冊用
CUresult cuLogsRegisterCallback(CUlogsCallback callbackFunc, void *userData,
                                CUlogsCallbackHandle *callback_out);

CUresult cuLogsUnregisterCallback(CUlogsCallbackHandle callback);

log 輸出管理 — 核心概念是 log iterator，它指向 buffer 的「目前結尾」：

// 取得目前 iterator；flags 目前必須為 0（其餘保留給未來版本）
CUresult cuLogsCurrent(CUlogIterator *iterator_out, unsigned int flags);

// 把 log buffer dump 到檔案或記憶體
CUresult cuLogsDumpToFile(CUlogIterator *iterator, const char *pathToFile,
                          unsigned int flags);
CUresult cuLogsDumpToMemory(CUlogIterator *iterator, char *buffer, size_t *size,
                            unsigned int flags);

dump 行為要點：

• iterator == NULL：dump 整個 buffer，最多 100 筆。
• iterator != NULL：從該筆開始 dump，並把 iterator 更新到目前 log 結尾（如同呼叫 cuLogsCurrent）。呼叫端可保存 iterator 位置，避免每次都 dump 整個 buffer。
• 若期間已寫入超過 100 筆，dump 開頭會加一行註記說明此情況（rollover）。
• 所有函式的 flags 目前都必須為 0，其餘選項保留給未來版本。

log buffer（環狀，最多 100 筆）
┌───┬───┬───┬─ ... ─┬────┐
│ 0 │ 1 │ 2 │       │ 99 │
└───┴───┴───┴─ ... ─┴────┘
              ▲iterator（指向目前結尾）
cuLogsCurrent → 取得 iterator
cuLogsDumpToFile/Memory(iterator) → 從 iterator 起 dump，並把 iterator 推到新結尾

cuLogsDumpToMemory 的額外注意事項：

1. buffer 整體會被 null-terminated，但個別 log entry 之間只用換行字元 (\n) 分隔。
2. buffer 最大尺寸為 25600 bytes。
3. 若 size 提供的空間不足以容納所有想要的 logs，會在第一筆加上註記，且最舊、塞不下的 entry 不會被 dump。
4. 函式返回後，size 會包含實際寫入 buffer 的位元組數。

Limitations and Known Issues（限制與已知問題）

1. log buffer 上限 100 筆；達上限後最舊的 entry 被覆寫，log dump 會包含一行 rollover 註記。
2. 尚未涵蓋所有 CUDA API，這是持續進行中的專案。
3. Error Log 的 log 位置（若有給定）在「真正產生 log」前不會被檢查有效性。
4. 這些 API 目前僅在 CUDA Driver 提供；未來版本會在 CUDA Runtime 加入對等 API。
5. log 訊息未在地化，所有 log 一律為美式英文 (US English)。

考試/測驗重點

Related Notes

• 04-CUDA-Features/01-Unified-Memory-Full-Support
• 04-CUDA-Features/06-Stream-Ordered-Memory-Allocator
• 04-CUDA-Features/25-Driver-Entry-Point-Access
• 01-Introduction-to-CUDA/05-CUDA-Platform
• 03-Advanced-CUDA/08-CUDA-Driver-API
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 04-CUDA-Features/Practice-CUDA-Features
• 00-Dashboard/Exam-Traps