錯誤檢查與函式/變數修飾符 (Error Checking and Specifiers)

重點總覽

2.1.7 Error Checking in CUDA

每個 CUDA API 都回傳一個列舉型別 cudaError_t。範例程式常省略檢查，但正式（production）應用應每次都檢查並管理回傳值。無錯誤時回傳 cudaSuccess。

許多應用會實作一個 utility macro 來統一檢查：

#define CUDA_CHECK(expr_to_check) do {                       \
    cudaError_t result = expr_to_check;                      \
    if (result != cudaSuccess) {                             \
        fprintf(stderr,                                      \
            "CUDA Runtime Error: %s:%i:%d = %s\n",           \
            __FILE__, __LINE__, result,                      \
            cudaGetErrorString(result));                     \
    }                                                        \
} while (0)

此 macro 使用 cudaGetErrorString API，把一個 cudaError_t 值轉成人類可讀字串。使用方式是把 API 呼叫包進 macro：

CUDA_CHECK(cudaMalloc(&devA, vectorLength * sizeof(float)));
CUDA_CHECK(cudaMalloc(&devB, vectorLength * sizeof(float)));

任一呼叫偵測到錯誤就會印到 stderr。小專案常用此 macro，大型應用可改接到 logging 系統。

• 關鍵事實
  • cudaError_t 是所有 runtime API 的統一回傳型別；cudaSuccess 代表成功。
  • cudaGetErrorString(result) 取得錯誤的文字描述。
  • macro 寫法可彈性改接 logging / 其他錯誤處理機制。

2.1.7.1 Error State（錯誤狀態）

CUDA runtime 為每個 host thread 維護一份 cudaError_t 狀態，預設 cudaSuccess，發生錯誤時被覆寫。

2.1.7.2 Synchronous vs Asynchronous Errors（同步與非同步錯誤）

CUDA kernel 啟動與許多 runtime API 都是非同步的。錯誤狀態在錯誤發生時被設定/覆寫，這代表非同步操作中發生的錯誤，只有在下次檢查錯誤狀態時才會被回報（可能是 cudaGetLastError、cudaPeekAtLastError，或任何回傳 cudaError_t 的 API）。

vecAdd<<<blocks, threads>>>(devA, devB, devC);
// 啟動後立即檢查錯誤狀態（偵測 launch 即時錯誤）
CUDA_CHECK(cudaGetLastError());
// 等待 kernel 執行完成；CUDA_CHECK 會回報執行期間發生的錯誤
CUDA_CHECK(cudaDeviceSynchronize());

兩段檢查的分工：cudaGetLastError() 抓 launch 當下的錯誤；cudaDeviceSynchronize() 等待執行並抓執行期間才發生的非同步錯誤。

非同步錯誤傳播時間軸
─────────────────────────────────────────────►
launch ── kernel 執行中 ──[非法存取!]── 之後的 API
                              │
                              └─ 錯誤狀態被設定，但 launch 當下檢查可能仍是 cudaSuccess
                                 之後每個回傳 cudaError_t 的 API 都會回傳此錯誤
                                 直到 cudaGetLastError() 清除狀態

2.1.7.3 CUDA_LOG_FILE

另一個辨識 CUDA 錯誤的好方法是 CUDA_LOG_FILE 環境變數。設定後，CUDA driver 會把遇到的錯誤訊息寫到該變數指定路徑的檔案。即使應用本身沒做回傳值檢查，也能擷取並辨識錯誤。

__global__ void k() { }
int main() {
    k<<<8192, 4096>>>();          // 非法 block 大小（4096 超過上限）
    CUDA_CHECK(cudaGetLastError());
    return 0;
}

$ env CUDA_LOG_FILE=cudaLog.txt ./errlog
CUDA Runtime Error: .../errorLogIllustration.cu:24:1 = invalid argument
$ cat cudaLog.txt
[CUDA][E] One or more of block dimensions of (4096,1,1) exceeds
         corresponding maximum value of (1024,1024,64)
[CUDA][E] Returning 1 (CUDA_ERROR_INVALID_VALUE) from cuLaunchKernel

• log 檔比 stderr 提供更多細節（哪個維度超過哪個上限、底層 cuLaunchKernel 回傳值）。
• 設為 stdout 或 stderr 會分別印到標準輸出/標準錯誤。

2.1.8 Device and Host Functions（函式修飾符）

• 函式（含 class member function、functor、lambda）可同時標註 __device__ 與 __host__，讓同一份原始碼在 CPU 與 GPU 兩端都能呼叫。

__host__ __device__ float square(float x) { return x * x; }   // 兩端皆可用

2.1.9 Variable Specifiers（變數修飾符）

CUDA 修飾符可用於靜態變數宣告以控制其放置位置。

2.1.9.1 Detecting Device Compilation（偵測 device 編譯）

當函式標 __host__ __device__，編譯器會同時產生 GPU 與 CPU 兩份程式碼。若想在其中只為 GPU 或只為 CPU 撰寫特定程式碼，最常見做法是用前處理器檢查是否定義了 __CUDA_ARCH__。

__host__ __device__ int f() {
#ifdef __CUDA_ARCH__
    // 只會編進 GPU 版本（device 編譯路徑）
    return device_path();
#else
    // 只會編進 CPU 版本（host 編譯路徑）
    return host_path();
#endif
}

• __CUDA_ARCH__ 只在 device 編譯階段有定義，故可用它分流。

__host__ __device__ 函式編譯分流
              ┌─────────────────────────┐
   原始碼 ───►│  nvcc                    │
              ├─────────────┬───────────┤
              │ device pass │ host pass │
              │ __CUDA_ARCH_│  未定義    │
              │ _ 有定義     │            │
              └─────┬───────┴─────┬─────┘
                    ▼             ▼
                 GPU 程式碼      CPU 程式碼

2.1.10 Thread Block Clusters（執行緒區塊叢集）

從 compute capability 9.0 起，CUDA programming model 加入一個可選的階層：thread block cluster，由多個 thread block 組成。

• 如同 thread block 內的 thread 保證同排程於同一個 SM，cluster 內的 thread block 也保證同排程於 GPU 內的同一個 GPC（GPU Processing Cluster）。
• cluster 同樣可組成 1D / 2D / 3D 的 cluster grid。

階層對應關係
 thread  ──(co-scheduled)──► SM   （在 thread block 內）
 block   ──(co-scheduled)──► GPC  （在 cluster 內）

 grid ⊃ cluster ⊃ block ⊃ thread

• cluster 大小：可由使用者定義；portable 上限為 8 個 block。硬體或 MIG 配置太小（不足 8 個 multiprocessor）時上限會相應下降。較小/較大（>8）配置為 architecture-specific，可用 cudaOccupancyMaxPotentialClusterSize API 查詢。
• 硬體同步：cluster 內所有 block 保證同時排程於單一 GPC，可用 cooperative groups API cluster.sync() 做硬體支援的同步。
• 查詢 API：cluster group 提供 num_threads()、num_blocks() 查 cluster 大小；dim_threads()、dim_blocks() 查 thread/block 在 cluster 內的 rank。
• Distributed Shared Memory：cluster 內 block 可存取 distributed shared memory（cluster 內所有 block shared memory 的合集），可對其中任意位址做讀、寫、atomics。

2.1.10.1 在 Triple Chevron 啟動 Cluster

啟用 cluster 有兩種方式：

// 編譯期 cluster 大小：X 維 2、Y/Z 維各 1
__global__ void __cluster_dims__(2, 1, 1)
cluster_kernel(float *input, float *output) { }

int main() {
    float *input, *output;
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    // grid 維度不受 cluster 影響，仍以 block 數計
    cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

考試/測驗重點

Related Notes

• 02-Programming-GPUs/01-CUDA-Cpp-Kernels-and-Launch
• 02-Programming-GPUs/02-CUDA-Cpp-Memory-Management
• 02-Programming-GPUs/03-CUDA-Cpp-Sync-and-Workflow
• 02-Programming-GPUs/05-CUDA-Python
• 02-Programming-GPUs/06-SIMT-Basics-and-Thread-Hierarchy
• 02-Programming-GPUs/07-SIMT-Device-Memory-Spaces
• 02-Programming-GPUs/09-SIMT-Atomics-Cooperative-Occupancy
• 02-Programming-GPUs/14-Async-Streams-and-Events
• 02-Programming-GPUs/16-Unified-and-System-Memory
• 02-Programming-GPUs/17-NVCC-Compiler
• 01-Introduction-to-CUDA/02-Execution-Model-and-SIMT
• 01-Introduction-to-CUDA/04-GPU-Memory-Hierarchy
• 02-Programming-GPUs/Practice-Programming-GPUs
• 00-Dashboard/Exam-Traps