NVIDIA Compute Sanitizer (NCS) 是一個功能強大的工具，它可以幫助您節省時間和精力，同時提高 CUDA 應用程序的可靠性和性能。
在我們之前的帖子 高效的 CUDA 調試：如何使用 NVIDIA Compute Sanitizer 追蹤 Bug 中，我們探討了并行編程領域的高效調試。我們討論了在 CUDA 環境中調試代碼的挑戰性和耗時性，尤其是在處理數千個線程時，以及 NCS 如何幫助實現這一過程。
這篇文章繼續我們對高效 CUDA 調試的探索。它重點介紹了更多的 NCS 工具，并介紹了幾個示例。
NVIDIA Compute Sanitizer
NCS 是一套工具，可以對代碼的功能正確性執行不同類型的檢查。NCS 中有四個主要工具：
- Memcheck 用于內存訪問錯誤和泄漏檢測
- Racecheck：這是一個用于檢測共享內存數據訪問風險的工具。
- Initcheck，這是一個用于檢測未初始化的設備全局內存訪問的工具
- Synccheck：用于線程同步的危險檢測。
除了這些工具之外，NCS 功能還包括：
- API，用于創建針對 CUDA 應用程序的清理和跟蹤工具
- 與 NVIDIA Tools 擴展集成（NVTX）
- CoreDump 支持 用于 CUDA-GDB
- 用于管理工具輸出的抑制功能
本文將重點介紹如何使用 initcheck 調試代碼和捕捉與未初始化的設備陣列相關的錯誤，以及如何使用同步檢查。請參閱 高效的 CUDA 調試：如何使用 NVIDIA Compute Sanitizer 追蹤 Bug 了解更多關于如何使用 memcheck 發現內存泄漏和跑道檢查查找競態條件的詳細信息。
初始化檢查
NCS Initcheck 幫助開發人員識別和解決 CUDA 代碼中未初始化的內存訪問錯誤。在 CUDA 應用程序中，未初始化的內存訪問可能導致不可預測的行為和不正確的結果。
NCS Initcheck 可以檢測設備代碼中對全局內存的未初始化內存訪問，并提供有關訪問位置和時間的詳細信息，以及訪問線程的堆棧跟蹤。這有助于揭示問題的根本原因并解決問題。
為了提供一個示例，下面的代碼受益于初始化檢查。
#include <stdio.h>#define THREADS 32#define BLOCKS 2__global__ void addToVector(float *v) { int tx = threadIdx.x + blockDim.x * blockIdx.x; v[tx] += tx;}int main(int argc, char **argv) { float *d_vec = NULL; float *h_vec = NULL; h_vec = (float *)malloc(BLOCKS*THREADS * sizeof(float)); cudaMalloc((void**)&d_vec, sizeof(float) * BLOCKS * THREADS); cudaMemset(d_vec, 0, BLOCKS * THREADS); // Zero the array addToVector<<<BLOCKS, THREADS>>>(d_vec); cudaMemcpy(h_vec, d_vec, BLOCKS*THREADS * sizeof(float), cudaMemcpyDeviceToHost); cudaDeviceSynchronize(); printf("After : Vector 0, 1 .. N-1: %f %f .. %f\n", h_vec[0], h_vec[1], h_vec[BLOCKS*THREADS-1]); cudaFree(d_vec); free(h_vec); exit(0);} |
此代碼包含一個名為addToVector它對向量中的每個元素執行簡單的值相加，并將結果寫回同一個元素。在?乍一看，它看起來很好：用庫達馬洛克，然后將其歸零cudaMemset，然后在內核中執行計算。它甚至打印出正確的答案：
$ nvcc -lineinfo initcheck_example.cu -o initcheck_example$ ./initcheck_exampleAfter : Vector 0, 1 .. N-1: 0.000000 1.000000 .. 63.000000 |
但代碼中包含一個小錯誤。（如果你能發現的話，得 20 分。）
使用 NCSinitcheck工具，用于檢查對設備上全局內存中矢量的任何訪問是否試圖讀取未初始化的值。
$ compute-sanitizer --tool initcheck ./initcheck_example========= COMPUTE-SANITIZER========= Uninitialized __global__ memory read of size 4 bytes========= at 0x70 in /home/pgraham/Code/BlogExamples/initcheck_example.cu:8:addToVector(float *)========= by thread (16,0,0) in block (0,0,0). . .========= Uninitialized __global__ memory read of size 4 bytes========= at 0x70 in /home/pgraham/Code/BlogExamples/initcheck_example.cu:8:addToVector(float *)========= by thread (17,0,0) in block (0,0,0). . . =========After : Vector 0, 1 .. N-1: 0.000000 1.000000 .. 63.000000========= ERROR SUMMARY: 48 errors |
這應該會打印很多信息（為了簡潔起見，顯示的輸出經過了編輯），但有些地方不正確。大量的輸出是回溯信息，可以使用–顯示回溯編號選項：
$ compute-sanitizer --tool initcheck --show-backtrace no ./initcheck_example |
查看輸出，您可以看到總共 48 個錯誤。報告顯示，它們都是這種類型，未初始化的__global__內存讀取大小為 4 字節.
每條消息都表示試圖從全局設備內存中讀取一些內容，并且這些內容的大小為 4 字節。一個合理的猜測是，錯誤指的是試圖訪問向量的元素，這些元素由每個 4 字節的浮點組成。
查看第一個錯誤，消息的下一部分指示是哪個線程和哪個線程塊導致了錯誤。在這種情況下，它是塊 0 中的線程 16。由于內核被設置為使得每個線程訪問向量的不同元素、向量的元素 17，d_vec[16]，未初始化。
在您的輸出中，您可能會看到一個不同的線程第一一個導致錯誤。GPU 可以按其認為合適的任何順序調度扭曲（32 個線程的組）。但是，檢查輸出的其余部分，并說服自己向量中導致錯誤的最低元素是元素 17（塊 0 中的線程 16）。
接下來，查看初始化（或應該初始化）數組的代碼行：
cudaMemset(d_vec, 0, BLOCKS * THREADS); // Zero the array |
檢查的定義cudaMemset，它需要三個參數：指向要設置的設備內存的指針(d_vec在這種情況下）每個字節在該內存區域中應該設置（在這種情況下為 0），以及要設置的字節數(塊*螺紋)。
現在問題開始變得更加明顯。矢量包含 64 個元素，由塊*螺紋構成，但每個元素都是一個浮點值，因此整個向量的長度為 256 字節。cudaMemset 僅初始化前 64 個 字節 （前 16 個元素），這意味著剩余的 192 個字節（相當于 48 個元素）未初始化。這 48 個元素對應于 48 個錯誤。
這與元素 17（線程 16，塊 0）是第一個導致錯誤的觀察結果相一致。賓果，發現問題了。
若要解決此問題，請更改cudaMemset呼叫：
cudaMemset(d_vec, 0, sizeof(float) * BLOCKS * THREADS); |
并檢查以確保消毒液是愉快的。
檢查未使用的內存
的另一個功能initcheck該工具正在識別應用程序結束時尚未訪問的已分配設備內存。在某些程序中，這可能是經過深思熟慮的——例如，使用大的靜態緩沖區來處理一系列潛在的問題大小。但是，當這更可能是一個導致錯誤的錯誤時，使用initcheck，如下所示。
#include <stdio.h>#define N 10__global__ void initArray(float* array, float value) { int threadGlobalID = threadIdx.x + blockIdx.x * blockDim.x; if (threadGlobalID < N) array[threadGlobalID] = value; return;}int main() { float* array; const int numThreadsPerBlock = 4; const int numBlocks = 2; cudaMalloc((void**)&array, sizeof(float) * N); initArray<<<numBlocks, numThreadsPerBlock>>>(array, 3.0); cudaDeviceSynchronize(); cudaFree(array); exit(0);} |
這個非常基本的代碼將揭示潛在的錯誤。它正在初始化一個數組，但線程的數量和塊的數量是硬編碼的。執行配置<<<…>>將啟動一個由八個線程組成的網格，而數據集有 10 個元素（最后兩個元素將不使用）。
使用 track unused memory（跟蹤未使用的內存）選項進行檢查。請注意，所需的語法將取決于所使用的 CUDA 版本。對于 12.3 之前的版本，使用以下內容提供參數“yes”：
--track-unused-memory yes ; |
從 12.3 版本開始，不需要提供參數，如下所示：
$ nvcc -o unused -lineinfo unused.cu$ compute-sanitizer --tool initcheck --track-unused-memory ./unused========= COMPUTE-SANITIZER========= Unused memory in allocation 0x7fe0a7200000 of size 40 bytes========= Not written 8 bytes at offset 0x20 (0x7fe0a7200020)========= 20% of allocation were unused.================== ERROR SUMMARY: 1 error |
清晰地跟蹤未使用的內存表示 40 個字節（10 x 4 字節浮動）的數組包含 8 個未寫入的字節。請使用數組地址（第一個長 0x…數字）和偏移量（0 x 20，十進制為 32，因此為 32 個字節或 8 個浮動）查看哪些字節未使用。正如預期的那樣，陣列中的浮動 9 和 10 沒有被使用。
要解決此問題，請使用N定義numBlocks:
const int numBlocks = (N + numThreadsPerBlock - 1) / numThreadsPerBlock; |
請注意–跟蹤未使用的內存設計用于分配的設備內存庫達馬洛克。該功能不適用于統一內存(cudaMallocManaged分配的存儲器）。
同步檢查
協作組編程模型啟用了在各種級別（不僅僅是塊和扭曲）同步線程的強大 CUDA 功能。協作組是一個設備代碼 API，用于定義、分區和同步線程組，相比標準提供了更多的靈活性和控制同步線程函數，用于同步塊中的所有線程。有關更多詳細信息，請參閱 協作組：靈活的 CUDA 線程編程。
然而，這種能力帶來了更多引入錯誤的機會。這就是 NCS同步檢查可以幫助識別和解決 CUDA 代碼中的同步錯誤。同步檢查可以識別 CUDA 應用程序是否正確地使用同步原語及其協同組 API 對應方。
同步的一個有趣的應用是將掩碼應用于線程的扭曲。設置扭曲，使一些線程為 true，另一些為 false，從而使每個線程能夠根據該屬性執行不同的操作。有關更多詳細信息，請參閱 使用 CUDA 扭曲級別基本體。
一個有用的功能是__氣球同步定義為：
unsigned int __ballot_sync(unsigned int mask, int predicate); |
面具是初始掩碼，通常在所有位都設置為 1 的情況下創建，表示扭曲中的所有線程最初都處于活動狀態。謂語是由每個線程計算的條件，其中謂詞對每個線程的計算結果為 true（非零）或 false（零）。
投票函數評估 warp 中每個線程的謂詞，并返回一個表示該線程結果的掩碼。它還提供了一個同步點。經線中的所有線都必須達到__氣球同步在他們中的任何一個能夠進一步進行之前。
例如，設置一個遮罩，其中扭曲中的偶數線程為 true，奇數線程為 false：
__ballot_sync(0xffffffff, threadID % 2 == 0); |
初始掩碼0xffffff是十六進制表示，計算結果為1111111111111111二進制。這確保了所有 32 個線程都參與到投票中。
投票結果是一個面具，0xaaaaaaaa，二進制形式為10101010101010101010偶數線程（線程 ID 0、2、4…）被設置為 true，奇數線程被設置為 false。
選票通常與__同步扭曲，可以基于所提供的掩碼同步經線中的線程。
以下示例同時使用_氣球同步和_同步扭曲:
static constexpr int NumThreads = 32 ;__shared__ int smem[NumThreads];__global__ void sumValues(int *sum_out) { int threadID = threadIdx.x; unsigned int mask = __ballot_sync(0xffffffff, threadID < (NumThreads / 2)); if (threadId <= (NumThreads / 2)) { smem[threadId] = threadId; __syncwarp(mask); if (threadID == 0) { *sum_out = 0; for (int i = 0; i < (NumThreads / 2); ++i) *sum_out += smem[i]; } } __syncThreads();}int main(){ int *sum_out = nullptr; cudaMallocManaged((void**)&sum_out, sizeof(int)); sumVaules<<<1, NumThreads>>>(sum_out); cudaDeviceSynchronize(); printf("Sum out = %d\n", *sum_out); cudaFree(sum_out); return 0;} |
在進一步閱讀之前，請先看一下代碼，并根據您對選票和同步扭曲功能。看看你是否能發現問題所在。（這次得了 50 分，更具挑戰性。）
這個代碼的目的是讓每個線程為共享內存分配一個值，然后將所有值相加得到一個答案。但是，這只適用于一半的可用線程。通過執行配置設置了 32 個線程的單個翹曲<<<1，numThreads>>執行內核sumValues。
在該內核中，使用__氣球同步具有threadID<線程數/2作為謂詞，它將在曲速的前半部分求值為 true，其中螺紋 ID<16（線程 0、1、..15）。
對于這 16 個線程，為共享內存分配一個值（threadID），然后執行__syncwarp（遮罩）同步這些線程以確保它們擁有所有? 寫入共享內存。然后基于這些值更新 sum_out 全局和。
接下來，嘗試編譯并運行以下代碼：
$ nvcc -o ballot_example -lineinfo ballot_example.cu$ ./ballot_exampleSum out = 0 |
答案為零，不正確。它應該是 120（15+14+13+…+2+1+0）。
你發現錯誤了嗎？代碼的條件部分使用 if 執行（threadId<=（線程數/2））。此代碼使用<=而不是<作為比較器，這意味著前 17 個線程將執行。
當線程 17 嘗試調用時會發生什么同步扭曲當它沒有被包含在面具中時？它? 導致整個內核停止運行，因此永遠不會達到總和計算。因此輸出為零。
所有這些都會無聲地失敗，只有不正確的輸出才表明有問題。在里面?更復雜的代碼，這可能是一場噩夢。
使用同步檢查提供了以下內容：
$ compute-sanitizer --tool synccheck --show-backtrace no ./ballot_example========= COMPUTE-SANITIZER========= Barrier error detected. Invalid arguments========= at 0x220 in /home/pgraham/Code/devblog/NCS_Part2/ballot_example.cu:32:sumValues(int *)========= by thread (0,0,0) in block (0,0,0)=========. . .========= Barrier error detected. Invalid arguments========= at 0x220 in /home/pgraham/Code/devblog/NCS_Part2/ballot_example.cu:32:sumValues(int *)========= by thread (16,0,0) in block (0,0,0)=========Sum out = 0========= ERROR SUMMARY: 17 errors |
關于這 17 個錯誤，“無效參數”synccheck 文檔 聲明，如果不是所有線程都到達 __同步扭曲，那么它們會在 mask 參數中聲明自己。
在這種情況下，線程 17 或線程（16,0,0）不是 活躍在掩碼中，所以它不應該調用同步扭曲。請注意，這會導致所有其他線程調用同步扭曲也登記一個錯誤。他們單獨打電話同步扭曲，但因為其中一個導致它失敗，而所有其他同步扭曲調用也必須失敗。這是一個集體操作，總共導致 17 個錯誤。
結論
這篇文章介紹了如何使用 NVIDIA Compute Sanitizer 中的 initcheck 和同步檢查功能。要開始使用 NCS，請下載 CUDA 工具包。
要了解更多信息，請訪問 NVIDIA/compute-sanitizer-samples 在 GitHub 上，并閱讀 NCS 文件。歡迎加入 NVIDIA 開發者論壇，這是一個專門討論 sanitize工具的平臺。祝您好運！
?
