CUDA 裝置端 Runtime (CUDA Device Runtime)

重點總覽

CUDA device runtime 是一套可在 kernel 內呼叫的 API，提供與 host 端 CUDA Runtime API 大致相同的能力，最常用於 Dynamic Parallelism (CDP2) 與 Device Graph Launch。原型在編譯時自動引入，不需手動 include cuda_device_runtime_api.h。它只是 host runtime 的一個子集：無多 GPU、無 cudaStreamSynchronize/cudaStreamQuery（CDP2 亦無 cudaDeviceSynchronize，可用 tail launch 取代其效果）、stream 必須以 cudaStreamNonBlocking 建立，並提供 fire-and-forget 與 tail launch 等特殊 named stream 取代同步。

引入與記憶體配置

引入 API

device runtime API 的原型在編譯時自動引入，與 host runtime 一樣，無需手動 #include <cuda_device_runtime_api.h>。

資源上限 (Configuration Options)

device runtime 系統軟體的資源由 host 程式透過 cudaDeviceSetLimit() 控制。

配置與生命週期 (Allocation and Lifetime)

cudaMalloc() / cudaFree() 在 host 與 device 環境語意不同：

• host 端 cudaMalloc()：從未使用的 device memory 配置新區段。
• device 端：映射到 device 側的 malloc() / free()，故總可配置量受 malloc heap size 限制，可能小於可用 device memory。

記憶體宣告

__device__ 與 __constant__

file scope 以 __device__ / __constant__ 宣告的變數，在 device runtime 下行為相同：

• 所有 kernel（無論由 host 或 device 啟動）皆可讀寫 __device__ 變數。
• 所有 kernel 對 module scope 的 __constant__ 都有相同視角。

Texture 與 Surface

• device code 不可建立或銷毀 texture / surface 物件。
• host 建立的 texture/surface 物件可在 device 自由使用與傳遞；動態建立的 texture 物件永遠有效，可由 parent 傳給 child kernel。

Shared Memory 宣告

兩種 shared memory 宣告在 device runtime 下都有效：靜態大小（file/function scope）與 extern 動態大小（由 launch 配置決定）。

__global__ void permute(int n, int *data) {
  extern __shared__ int smem[];
  if (n <= 1)
    return;
  smem[threadIdx.x] = data[threadIdx.x];
  __syncthreads();
  permute_data(smem, n);
  __syncthreads();
  // SMEM 不能傳給 child，故寫回 GMEM
  data[threadIdx.x] = smem[threadIdx.x];
  __syncthreads();

  if (threadIdx.x == 0) {
    permute<<< 1, 256, n/2*sizeof(int) >>>(n/2, data);
    permute<<< 1, 256, n/2*sizeof(int) >>>(n/2, data+n/2);
  }
}

void host_launch(int *data) {
  permute<<< 1, 256, 256*sizeof(int) >>>(256, data);
}

Constant Memory 與 Symbol Addresses

• Constant 不可從 device 修改；只能從 host 改。若 host 在某 grid 存取該 constant 的生命週期內並行修改它，行為未定義。
• __device__ symbol 可在 kernel 內直接用 & 取址（皆在可見位址空間）；__constant__ symbol 也可取址，但指標指向唯讀資料。
• 因可直接引用 symbol，cudaMemcpyToSymbol()、cudaGetSymbolAddress() 這類引用 symbol 的 API 不需要也不支援。故執行中的 kernel 即使在 child launch 前也不能改 constant data。

Launch 機制與 Device 管理

SM Id 與 Warp Id

PTX 的 %smid 與 %warpid 是 volatile 值。device runtime 可能把 thread block 重新排到不同 SM 以更有效管理資源。

Launch Setup APIs

device 端 kernel launch 使用與 host 相同的 triple chevron <<<>>> 語法。底層由 device runtime library 的系統級機制實現，也可從 PTX 直接以 cudaGetParameterBuffer() 與 cudaLaunchDevice() 呼叫。NVCC front-end 把 <<<>>> 翻譯成這些呼叫。允許應用程式自行呼叫這兩個 API（要求與 PTX 相同），但需自行依規格正確填妥所有必要資料結構；CUDA 保證這些資料結構向後相容（backwards compatibility）。

extern __device__ cudaError_t cudaGetParameterBuffer(void **params);
extern __device__ cudaError_t cudaLaunchDevice(void *kernel,
                                               void *params, dim3 gridDim,
                                               dim3 blockDim,
                                               unsigned int sharedMemSize = 0,
                                               cudaStream_t stream = 0);

Device 管理

• device runtime 無多 GPU 支援：只能操作當前正在執行的 device。
• 但允許查詢系統中任意 CUDA capable device 的屬性。

支援的 API 子集

host 與 device runtime API 語法相同；語意除非特別標註否則一致。

API 錯誤與 Launch 失敗

• 任何函式都可能回傳錯誤碼（型別 cudaError_t）；最後一個錯誤碼以 per-thread 記錄，用 cudaGetLastError() 取得。
• device 端 launch 可能因多種原因失敗；需呼叫 cudaGetLastError() 判斷 launch 是否出錯。

ECC Errors

kernel 內的 code 無法收到 ECC 錯誤通知。ECC 錯誤在整個 launch tree 完成後於 host 側回報；nested 程式執行中發生的 ECC 錯誤會產生例外或繼續執行（視錯誤與配置而定）。

Device Runtime Streams

device runtime 提供 named 與 unnamed (NULL) stream。

Implicit (NULL) Stream

device runtime 提供單一 implicit 無名 stream，由 thread block 內所有 thread 共享。但因所有 named stream 都以 cudaStreamNonBlocking 建立，NULL stream 的工作不會對其他 stream（含其他 thread block 的 NULL stream）的待處理工作插入隱式相依——即 host 端 NULL stream 的 cross-stream barrier 語意在 device 不支援。

Fire-and-Forget Stream

cudaStreamFireAndForget：以更少 boilerplate、無 stream 追蹤開銷送出 fire-and-forget 工作；功能等同於「每次 launch 建新 stream」但更快。

• 立即排程，不相依於先前 grid 的完成。
• 沒有其他 grid launch 能相依於 fire-and-forget 的完成（除了 parent grid 結尾的隱式同步）——故 tail launch 或 parent stream 的下一個 grid 都不會在 parent 的 fire-and-forget 工作完成前啟動。

// C2 的 launch 不會等 C1 完成
__global__ void P( ... ) {
  C1<<< ... , cudaStreamFireAndForget >>>( ... );
  C2<<< ... , cudaStreamFireAndForget >>>( ... );
}

Tail Launch Stream

cudaStreamTailLaunch：讓一個 grid 排程「在自己完成之後才啟動」的新 grid，多數情況可達成等同 cudaDeviceSynchronize() 的效果。

• 每個 grid 有自己的 tail launch stream。
• parent 的所有 非 tail 工作（ordinary / per-thread / fire-and-forget）會在 tail stream 啟動前隱式同步完成。
• 送入同一 grid 的 tail launch stream 的兩個 grid 依序執行：後者要等前者及其所有後代完成。

// C2 只在 C1 完成後 launch
__global__ void P( ... ) {
  C1<<< ... , cudaStreamTailLaunch >>>( ... );
  C2<<< ... , cudaStreamTailLaunch >>>( ... );
}

// C 只在所有 X、F、P 完成後 launch
__global__ void P( ... ) {
  C<<< ... , cudaStreamTailLaunch >>>( ... );
  X<<< ... , cudaStreamPerThread >>>( ... );
  F<<< ... , cudaStreamFireAndForget >>>( ... );
}

tail launch stream 的行為「如同被插入在 parent grid 與 parent stream 的下一個 grid 之間」——parent stream 的下一個 grid 不會在 parent 的 tail 工作完成前啟動。

parent stream 順序:
  ... → [P1] → [P1 的 tail: C] → [P2] → ...
            └─ C 完成後 P2 才 launch

每個 grid 只有一條 tail launch stream；要 tail launch 並行多個 grid，可借一層 helper grid 用 fire-and-forget 展開：

// C1 與 C2 在 P 完成後並行 launch
__global__ void T( ... ) {
  C1<<< ... , cudaStreamFireAndForget >>>( ... );
  C2<<< ... , cudaStreamFireAndForget >>>( ... );
}
__global__ void P( ... ) {
  ...
  T<<< ... , cudaStreamTailLaunch >>>( ... );
}

考試/測驗重點

Related Notes

• 05-Technical-Appendices/12-Memory-Barrier-Pipeline-Cooperative-APIs
• 05-Technical-Appendices/02-CUDA-Environment-Variables
• 04-CUDA-Features/22-Dynamic-Parallelism
• 03-Advanced-CUDA/02-Advanced-Streams-and-Dependent-Launch
• 05-Technical-Appendices/Practice-Technical-Appendices
• 00-Dashboard/Exam-Traps