? ? ? ? 調試 CUDA Kernel 并進入?__device__
?函數是 CUDA 開發中一項非常重要的技能。這主要依賴于 NVIDIA 的官方調試器?NVIDIA Nsight Systems?(用于系統級分析) 和?NVIDIA Nsight Compute?(用于內核級分析) 以及經典的?cuda-gdb?(命令行調試器)。
這里將重點介紹兩種最常用和強大的方法:使用?Nsight Visual Studio Code Edition?(圖形化界面,推薦) 和?cuda-gdb?(命令行)。
?
方法一:使用 Nsight VSCode Edition (圖形化界面,最推薦)
? ? ? ? Nsight for VSCode 提供了類似于調試 CPU 代碼的直觀體驗,是當前調試 CUDA 的首選工具。
前提條件
-
安裝 NVIDIA 驅動和 CUDA Toolkit: 確保你的系統安裝了正確版本的驅動和 CUDA。
-
安裝 VSCode: 從?VSCode 官網?下載并安裝。
-
安裝 Nsight VSCode 插件: 在 VSCode 的擴展市場中搜索 “NVIDIA Nsight” 并安裝。
-
確保你的代碼是可調試的: 在編譯你的 CUDA 代碼時,必須使用?
-G
?或?-lineinfo
?標志來生成調試信息。-
-G
: 生成完整的調試信息,但會嚴重禁用所有編譯器優化,極大影響性能,僅用于調試。 -
-lineinfo
?(--generate-line-info
): 生成行號信息,允許調試和性能分析,但對性能影響較小,是?Nsight Compute
?性能分析的首選。對于調試?__device__
?函數,-lineinfo
?通常就足夠了。
示例編譯命令:
nvcc -g -G -o my_program my_program.cu # 使用 -g -G 進行完整調試 # 或者 nvcc --device-debug -o my_program my_program.cu # CUDA 11.2+ 推薦方式,等同于 -g -G # 或者 (如果主要為了分析,附帶一些調試能力) nvcc -lineinfo -o my_program my_program.cu
-
調試步驟
-
打開項目: 在 VSCode 中打開你的 CUDA 項目文件夾。
-
創建調試配置文件:
-
點擊側邊欄的 “運行和調試” 圖標 (或?
Ctrl+Shift+D
)。 -
點擊 “create a launch.json file”。
-
在彈出的環境選擇中,選擇?CUDA C++。
-
這會在項目下生成一個?
.vscode/launch.json
?文件。
-
-
配置?
launch.json
:-
一個基本的配置如下。關鍵是指定正確的程序路徑和?
cuda-gdb
?路徑。
{"version": "0.2.0","configurations": [{"name": "CUDA C++: Launch","type": "cuda-gdb","request": "launch","program": "${workspaceFolder}/my_program", // 你的可執行文件路徑"stopAtEntry": false, // 設為 true 會在 main 函數入口處暫停"cwd": "${workspaceFolder}","args": [], // 傳遞給程序的命令行參數"environment": [], // 環境變量"externalConsole": false}] }
-
-
設置斷點:
-
在你的 CUDA 源碼 (
.cu
?或?.cuh
?文件) 中,在你想要中斷的行的左側空白處點擊。 -
你可以在?
main
?函數、<<<...>>>
?調用的 kernel 函數、以及任何?__device__
?函數中設置斷點。
-
-
開始調試:
-
選擇剛剛創建的 “CUDA C++: Launch” 配置。
-
按?
F5
?或點擊綠色的 “開始調試” 按鈕。 -
程序開始運行,并在遇到你設置的斷點時暫停。
-
-
步入?
__device__
?函數:-
當程序在 kernel 的某一行暫停時,你可以使用調試控制欄的按鈕:
-
Step Into (
F11
): 如果當前行調用了某個?__device__
?函數,按?F11
?會進入該?__device__
?函數的函數體。 -
Step Over (
F10
): 執行當前行,但不進入函數內部。 -
Step Out (
Shift+F11
): 執行完當前函數,返回到調用它的地方。
-
-
-
查看變量和調用堆棧:
-
在調試過程中,你可以在 VSCode 的左側面板查看變量的當前值。
-
調用堆棧?面板會顯示你當前的執行位置,從?
main
?到 kernel 再到?__device__
?函數,清晰明了。 -
你還可以將鼠標懸停在源碼中的變量上來查看其值。
-
?
方法二:使用 cuda-gdb (命令行)
? ? ? ? 但其實我們更習慣命令行操作,或者在沒有圖形界面的遠程服務器上工作,cuda-gdb
?是強大的選擇。
前提條件
? ? ? ? ? 同樣,編譯時必須使用?-G
?或?-lineinfo
?標志。
nvcc -g -G -o my_program my_program.cu
例如:?nvcc -g -G --gpu-architecture=sm_120? sgemm_1.cu -o sgemm_1 -I ../cutlass/include/ -I ../cutlass/tools/util/include -I /usr/local/cuda/include
調試步驟
-
啟動調試器:
cuda-gdb ./my_program
-
設置斷點:
-
在 kernel 函數處設置斷點:
(cuda-gdb) break my_kernel_function
-
在?
__device__
?函數處設置斷點:(cuda-gdb) break my_device_function
-
在特定行設置斷點:
(cuda-gdb) break file.cu:123
-
-
運行程序:
(cuda-gdb) run
程序會開始執行,并在第一個斷點處停止。
-
控制執行和步入?
__device__
?函數:-
next
?(n
): Step Over,執行下一行。 -
step
?(s
):?Step Into,進入函數。如果下一行是?__device__
?函數調用,這會進入該?__device__
?函數體。 -
continue
?(c
): 繼續運行直到下一個斷點。 -
finish
: Step Out,運行到當前函數返回。
-
-
檢查線程和變量:
-
CUDA 調試的核心是理解線程。你可以切換當前關注的線程:
(cuda-gdb) cuda thread 1 # 切換到 block 0, thread 0 (cuda-gdb) cuda thread (1, 2, 3) # 切換到 blockIdx(1,2), threadIdx(3)
-
打印變量值:
(cuda-gdb) print variable_name
-
-
退出:
?(cuda-gdb) quit
重要提示和常見問題
-
硬件兼容性: 并非所有 NVIDIA GPU 都支持調試功能。請查閱?NVIDIA 文檔?確認你的 GPU 是否支持 “TCC 模式” 或 “Compute Preemption”,這對于調試至關重要。消費級顯卡 (GeForce) 的調試支持可能不如專業卡 (Tesla, Quadro) 完善。
-
性能影響: 使用?
-G
?編譯會極大降低 kernel 運行速度,并且調試本身也會引入開銷。這是正常的,目的是為了獲得精確的調試狀態。 -
焦點線程: 在 CUDA 調試中,任何時候你都只在一個特定的 GPU 線程上下文中查看變量和執行代碼。確保你正在檢查你感興趣的線程。
-
Nsight vs cuda-gdb:?
Nsight VSCode
?底層調用的也是?cuda-gdb
,但它提供了無比友好的圖形界面,極大地簡化了操作,強烈推薦初學者和絕大多數開發者使用。
總結:要調試并進入?__device__
?函數,只需?1) 用?-G
?編譯,2) 在 Nsight VSCode 或 cuda-gdb 中設置斷點,3) 使用?Step Into
?(F11 /?step
) 命令即可。