CUDA Driver API (The CUDA Driver API)

重點總覽

Driver API 概觀

CUDA runtime（即 cuda 前綴的 runtime API）是建構在更低階的 CUDA driver API 之上。driver API 實作於裝置驅動安裝時複製到系統的動態函式庫（cuda.dll 或 cuda.so），所有入口點皆以 cu 為前綴。

• 它是 handle-based、imperative（命令式） 的 API：大多數物件以不透明 handle 引用，傳給函式來操作物件。
• 相對 runtime API 提供更低階的明確控制：context、module 載入、kernel 啟動等都要手動完成。
• runtime 與 driver API 可互通（見下方 Interoperability）。

下表為 driver API 中常見物件與其 handle 型別：

整體使用流程：先 cuInit()，再建立並 current 一個 context；於 context 內以 host code 明確載入 PTX/binary module；再用 cuLaunchKernel() 等入口點啟動 kernel。

下面是 driver API 版 VecAdd 的 host code 骨架（注意 cuInit → cuDeviceGet → cuCtxCreate → cuModuleLoad → cuMemAlloc → cuMemcpyHtoD → cuModuleGetFunction → cuLaunchKernel 的順序）：

cuInit(0);                                   // 必須最先呼叫

int deviceCount = 0;
cuDeviceGetCount(&deviceCount);

CUdevice cuDevice;
cuDeviceGet(&cuDevice, 0);                    // device 0 的 handle

CUcontext cuContext;
cuCtxCreate(&cuContext, 0, cuDevice);         // 建立 context 並 current

CUmodule cuModule;
cuModuleLoad(&cuModule, "VecAdd.ptx");        // 載入 module

CUdeviceptr d_A, d_B, d_C;
cuMemAlloc(&d_A, size);
cuMemAlloc(&d_B, size);
cuMemAlloc(&d_C, size);
cuMemcpyHtoD(d_A, h_A, size);
cuMemcpyHtoD(d_B, h_B, size);

CUfunction vecAdd;
cuModuleGetFunction(&vecAdd, cuModule, "VecAdd");

void* args[] = { &d_A, &d_B, &d_C, &N };
cuLaunchKernel(vecAdd,
               blocksPerGrid, 1, 1,           // grid dims
               threadsPerBlock, 1, 1,         // block dims
               0, 0, args, 0);                // sharedMem, stream, params, extra

要點：runtime API 用一行 cudaLaunchKernel/<<<>>> 隱式完成的事，driver API 要逐步手動建立 context、載入 module、取得 function handle 再啟動。完整程式見 vectorAddDrv CUDA sample。

3.3.1 Context

一個 CUDA context 類比於一個 CPU process。driver API 內所有資源與動作都封裝在某個 context 內，當 context 被銷毀時系統自動清理這些資源。

• 除了 module、texture/surface reference 等物件外，每個 context 有自己獨立的 address space。因此不同 context 的 CUdeviceptr 指向不同的記憶體位置。
• 一個 host thread 同一時間只能有一個 device context 為 current。cuCtxCreate() 建立時即把該 context current 給呼叫的 host thread。
• 在 context 中運作的 CUDA 函式（大多數不涉及 device 列舉或 context 管理者），若無有效 current context，會回傳 CUDA_ERROR_INVALID_CONTEXT。

Context stack：每個 host thread 維護一個 current contexts 的堆疊。

• cuCtxCreate() 把新 context push 到堆疊頂端。
• cuCtxPopCurrent() 把 context 從 host thread detach；此時 context 變為「floating」，可被 push 成任一 host thread 的 current context。cuCtxPopCurrent() 同時還原前一個 current context（若有）。

host thread 的 context stack（堆疊頂 = current）

  cuCtxCreate / cuCtxPushCurrent        cuCtxPopCurrent
        push ↓                              pop ↑
   ┌──────────────┐  <- current(top)   ┌──────────────┐  <- floating（detach 後）
   │  ctxB (new)  │                    │  ctxB        │ ──→ 可 push 給其他 thread
   ├──────────────┤                    └──────────────┘
   │  ctxA        │                    ┌──────────────┐
   └──────────────┘                    │  ctxA        │  <- 還原成 current
                                       └──────────────┘

Usage count（使用計數）：每個 context 維護一個使用計數。

• cuCtxCreate() 建立時使用計數為 1。
• cuCtxAttach() 遞增、cuCtxDetach() 遞減使用計數。
• 當使用計數歸 0（透過 cuCtxDetach() 或 cuCtxDestroy()）時，context 被銷毀。
• usage count 便於同一 context 中由第三方撰寫之程式碼互通：例如三個函式庫共用一個 context，各自 cuCtxAttach() 增計數、用完 cuCtxDetach() 減計數。

3.3.2 Module

Module 是可動態載入的 device code 與資料套件，類比 Windows 的 DLL，由 nvcc 輸出（見 Compilation with NVCC）。

• 所有符號名稱（functions、global variables、texture/surface references）維護在 module scope，讓獨立第三方撰寫的 module 能在同一 CUDA context 內互通。
• 載入 module 並取得 kernel handle：

CUmodule cuModule;
cuModuleLoad(&cuModule, "myModule.ptx");      // 從檔案載入 cubin/PTX/fatbin
CUfunction myKernel;
cuModuleGetFunction(&myKernel, cuModule, "MyKernel");   // 依名稱取 kernel

從 PTX 即時編譯並解析錯誤：用 cuModuleLoadDataEx 搭配 CUjit_option 選項陣列傳入 error log buffer，可在 JIT 失敗時取得編譯錯誤訊息。

CUjit_option options[3];
void* values[3];
char error_log[BUFFER_SIZE];
options[0] = CU_JIT_ERROR_LOG_BUFFER;            values[0] = (void*)error_log;
options[1] = CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES; values[1] = (void*)BUFFER_SIZE;
options[2] = CU_JIT_TARGET_FROM_CUCONTEXT;       values[2] = 0;
err = cuModuleLoadDataEx(&cuModule, PTXCode, 3, options, values);
if (err != CUDA_SUCCESS) printf("Link error:\n%s\n", error_log);

多個 PTX 的編譯 + 連結：用 cuLinkCreate → cuLinkAddData（每段 PTX 一次，input type CU_JIT_INPUT_PTX）→ cuLinkComplete 產生 cubin → cuModuleLoadData 載入，最後 cuLinkDestroy。可透過 CU_JIT_WALL_TIME、CU_JIT_INFO_LOG_BUFFER、CU_JIT_LOG_VERBOSE 取得連結時間與資訊。

cuLinkCreate(6, options, values, &linkState);
cuLinkAddData(linkState, CU_JIT_INPUT_PTX, (void*)PTXCode0, strlen(PTXCode0)+1, 0,0,0,0);
cuLinkAddData(linkState, CU_JIT_INPUT_PTX, (void*)PTXCode1, strlen(PTXCode1)+1, 0,0,0,0);
cuLinkComplete(linkState, &cubin, &cubinSize);   // 產生連結後的 cubin
cuModuleLoadData(cuModule, cubin);
cuLinkDestroy(linkState);

加速載入：載入 cubin 等過程可用多執行緒加速，傳 CU_JIT_BINARY_LOADER_THREAD_COUNT（值設 0 表示使用機器上 CPU 數量的執行緒）。完整程式見 ptxjit CUDA sample。

3.3.3 Kernel Execution

cuLaunchKernel() 以給定的 execution configuration 啟動一個 kernel。參數有兩種傳遞方式：

對齊規則（用 buffer 傳參時關鍵）：

• built-in vector types 的 device 端對齊需求列於文件 Table 43。
• 其他基本型別：device 端對齊 = host 端對齊，可用 __alignof() 取得。
• 例外：當 host 編譯器把 double 與 long long（64-bit 系統上含 long）對齊在 one-word 邊界（如 gcc -mno-align-double），但 device code 中這些型別永遠對齊在 two-word 邊界。
• CUdeviceptr 是整數但代表指標，對齊需求為 __alignof(void*)。

用 macro 逐一對齊並填入 parameter buffer，再經 extra[] 傳入：

#define ALIGN_UP(offset, alignment) \
    (offset) = ((offset) + (alignment) - 1) & ~((alignment) - 1)

ADD_TO_PARAM_BUFFER(i,  __alignof(i));
ADD_TO_PARAM_BUFFER(f4, 16);            // float4 對齊為 16
ADD_TO_PARAM_BUFFER(devPtr, __alignof(devPtr));
ADD_TO_PARAM_BUFFER(f2, 8);             // float2 對齊為 8

void* extra[] = {
    CU_LAUNCH_PARAM_BUFFER_POINTER, paramBuffer,
    CU_LAUNCH_PARAM_BUFFER_SIZE,     &paramBufferSize,
    CU_LAUNCH_PARAM_END
};
cuLaunchKernel(cuFunction,
               blockWidth, blockHeight, blockDepth,
               gridWidth, gridHeight, gridDepth,
               0, 0, 0, extra);          // args 為 0，改用 extra 傳參

typedef struct {
    float  f;
    float4 f4;
} myStruct;

3.3.4 Interoperability between Runtime and Driver APIs

應用程式可將 runtime API 與 driver API 程式碼混用。

• 若 context 由 driver API 建立並 current，後續 runtime 呼叫會使用此 context，而不另建一個。
• 若 runtime 已初始化，可用 cuCtxGetCurrent() 取得初始化時建立的 context，供後續 driver API 使用。
• runtime 隱式建立的 context 稱為 primary context，可從 driver API 以 Primary Context Management 函式管理。
• device memory 可用任一 API 配置/釋放；CUdeviceptr 可與一般指標互轉：

CUdeviceptr devPtr;
float* d_data;
// 用 driver API 配置
cuMemAlloc(&devPtr, size);
d_data = (float*)devPtr;
// 用 runtime API 配置
cudaMalloc(&d_data, size);
devPtr = (CUdeviceptr)d_data;

• 因此 driver API 寫的應用程式可呼叫以 runtime API 撰寫的函式庫（如 cuFFT、cuBLAS）。
• reference manual 中 device management 與 version management 兩節的所有函式皆可互換使用。

context 來源與互通

 driver API ──cuCtxCreate──► 自建 context ──► 後續 runtime 呼叫沿用同一 context
 runtime   ──隱式初始化──► primary context ◄── cuCtxGetCurrent / cuDevicePrimaryCtxRetain
                                              （driver 端以 Primary Context Mgmt 管理）

 CUdeviceptr  ⇄  float* / void*   （可雙向 cast）

考試/測驗重點

Related Notes

• 03-Advanced-CUDA/01-Advanced-Launch-and-Clusters
• 03-Advanced-CUDA/03-Batched-Transfers-and-Env-Vars
• 03-Advanced-CUDA/04-Using-PTX-and-Hardware-Model
• 03-Advanced-CUDA/09-Multi-GPU-Programming
• 03-Advanced-CUDA/10-Tour-of-CUDA-Features
• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 02-Programming-GPUs/17-NVCC-Compiler
• 01-Introduction-to-CUDA/05-CUDA-Platform
• 03-Advanced-CUDA/Practice-Advanced-CUDA
• 00-Dashboard/Exam-Traps