cuda-memcheck : CUDA Debug 好幫手

在撰寫 CUDA 程式的時候最煩惱的就是 Debug 困難

即使現在有了 cuda-gdb 也是相當的困難

而在 CUDA toolkit 裡面有附贈一隻叫 cuda-memcheck 的小程式

可以幫助你檢查程式中有無下列兩種錯誤：


1. Out of boundary memory access
2. Misalignment

那接下來先了解一下上面兩個錯誤是啥

第一個錯誤 Out of boundary memory access 很好理解

比如說你的陣列長度只有 10 但是你卻存取到了 10 以上

int a[10];
a[11] = 1111; // Out of boundary memory access !!

這件事情在傳統 C 語言上面很好察覺,

但在 CUDA 中通常 index 都是跟 threadIdx 以及 blockIdx 加起來之類的

一不小心就會寫錯

而第二個錯誤 Misalignment 則是一般在 x86 PC 上寫程式幾乎不會遇到的事情

原因在於 x86 硬體允許你作這件事, 在嵌入式系統的世界比較容易遇到這種事

alignment 這件事簡單來講就是記憶體存取時位址必須為 x 的倍數

例如 4 - byte alignment 的話, 位址必須能被 4 整除

合法的 4 - byte alignment  位址如 : 0x00ffff, 0x0, 0x8 等等
不合法的 4 - byte alignment  位址如 : 0x01, 0xff06, 0xa 等等

但通常這種情只會發生在你有對指標作些特殊處理時才會發生

例如

char p[10];
int *a = (int*)p+1;

之類的運算, 通常在寫影像處理相關的程式時
比較有可能寫出這類的程式碼
在寫 CUDA 時請盡量避開

在了解 cuda-memchech 的功能後來看看如何使用

假設你有一個 CUDA 程式 a.out

那你只要下

cuda-memcheck a.out

就可以開始檢查

如果有錯誤的話會吐出類似下面的訊息


========= CUDA-MEMCHECK
========= Invalid __global__ read of size 4
=========     at 0x00000038 in f
=========     by thread (0,0,0) in block (0,0)
=========     Address 0xfb00000228 is out of bounds
=========
========= ERROR SUMMARY: 1 error


在這個時候你會發現雖然他告訴你 out of bounds
但是根本不知道在那一行出錯阿！！！

這時候你重新編譯你的程式

nvcc -G oob.cu

加個 -G 然後

cuda-memcheck a.out

重新檢查


========= CUDA-MEMCHECK
========= Invalid __global__ read of size 4
=========     at 0x000000b0 in oob.cu:3:f
=========     by thread (0,0,0) in block (0,0)
=========     Address 0xfb00000228 is out of bounds
=========
========= ERROR SUMMARY: 1 error

就會報告錯誤的行數了！

另外有個 --continue flag

cuda-memcheck --continue ./a.out

這時候他檢測到錯誤的時候就會試著繼續跑跑看了

Misalignment 錯誤訊息差不多也是這樣就不浪費版面了