本文翻譯自：https://developer.nvidia.com/blog/using-nsight-compute-to-inspect-your-kernels/
使用 Nsight Compute 檢查您的內核
Nvidia 為開發人員提供了新的 CUDA 工具庫：Nsight Compute 和 Nsight Systems 。在更新的 GPU 架構上，這些工具將會變得更重要。在本博客的示例中，我們將會使用這些新工具來獲取圖靈及更高架構的計算結果。建議先閱讀本系列的前兩篇博客，以獲取相關信息：
如前幾篇博客所述，Nsight Compute 和Nsight Systems 的目標和功能是不同的，所以調優行為會由一個或者幾個這種新工具組合使用。Nsight Compute 的主要用途之一是提供對 Kernel 的 GPU 性能分析指標。如果您使用過 NVIDIA Visual Profiler 或 nvprof（命令行分析器），您可能已經檢查了 CUDA 內核的特定指標。本博客重點介紹如何使用 Nsight Compute 做到這一點。許多其他的分析功能（例如檢查時間線、測量活動持續時間等）可以使用 Nsight Systems 執行。
開始
我們將分析一段 CUDA 代碼，它是上一篇博客中向量相加的代碼的變體。這段代碼使用了 2D CUDA grid 配置 ，以及二維數組(即雙下標)訪問。
#include
const size_t size_w = 1024;
const size_t size_h = 1024;
typedef unsigned mytype;
typedef mytype arr_t[size_w];
const mytype A_val = 1;
const mytype B_val = 2;
__global__ void matrix_add_2D(const arr_t * __restrict__ A, const arr_t * __restrict__ B, arr_t * __restrict__ C, const size_t sw, const size_t sh){
size_t idx = threadIdx.x+blockDim.x*(size_t)blockIdx.x;
size_t idy = threadIdx.y+blockDim.y*(size_t)blockIdx.y;
if ((idx < sh) && (idy < sw)) C[idx][idy] = A[idx][idy] + B[idx][idy];
}
int main(){
arr_t *A,*B,*C;
size_t ds = size_w*size_h*sizeof(mytype);
cudaError_t err = cudaMallocManaged(&A, ds);
if (err != cudaSuccess) {std::cout << "CUDA error: " << cudaGetErrorString(err) << std::endl; return 0;}
cudaMallocManaged(&B, ds);
cudaMallocManaged(&C, ds);
for (int x = 0; x < size_h; x++)
for (int y = 0; y < size_w; y++){
A[x][y] = A_val;
B[x][y] = B_val;
C[x][y] = 0;}
int attr = 0;
cudaDeviceGetAttribute(&attr, cudaDevAttrConcurrentManagedAccess,0);
if (attr){
cudaMemPrefetchAsync(A, ds, 0);
cudaMemPrefetchAsync(B, ds, 0);
cudaMemPrefetchAsync(C, ds, 0);}
dim3 threads(32,32);
dim3 blocks((size_w+threads.x-1)/threads.x, (size_h+threads.y-1)/threads.y);
matrix_add_2D<<<blocks,threads>>>(A,B,C, size_w, size_h);
cudaDeviceSynchronize();
err = cudaGetLastError();
if (err != cudaSuccess) {std::cout << "CUDA error: " << cudaGetErrorString(err) << std::endl; return 0;}
for (int x = 0; x < size_h; x++)
for (int y = 0; y < size_w; y++)
if (C[x][y] != A_val+B_val) {std::cout << "mismatch at: " << x << "," << y << " was: " << C[x][y] << " should be: " << A_val+B_val << std::endl; return 0;} ;
std::cout << "Success!" << std::endl;
return 0;
}
重點：
-
Managed Memory：我們使用Managed Memory 進行數據分配。對於支持頁面級別的缺頁異常的 GPU （即Unified Memory 2.0）我們預取數據以避免對內核造成性能影響。
-
2D：我們使用了 2D 的 grid 和 block 形狀，並且使用 typedef 來簡化 2D 數據的定義，其中數據寬度在編譯時是已知的（在此示例中）。這讓我們可以使用雙下標進行訪問，而不需要指針操作。
-
內核設計：內核非常簡單。每個線程使用CUDA內置變量計算一組 x、y 下標，如果計算的下標是有效的(在有效的數據區域內)，就將所選元素相加。
希望上面的代碼看起來很簡單。作為一名 CUDA 開發人員，你應該知道兩個最重要的優化事項：給 GPU 足夠多的並行任務、有效使用顯存子系統。我們將着重於第二個目標。因為我們僅使用了全局存儲器，因此我們的重點是如何有效的使用全局存儲器，為此我們需要盡量對全局存儲器進行讀/寫聯合操作。
在 Visual Profiler (nvvp) 或 nvprof 中，可以使用 gld_efficiency （全局讀取效率）和 gst_efficiency （全局寫入效率）等指標，可以輕松的驗證全局存儲器的訪問是否已經合並。
有哪些指標？
一般來說，Nsight Compute 所使用的指標與以往的工具不同。例如，目前 Nsight Compute 還沒有提供與以前 gld_efficiency 和 gst_efficiency 相對應的指標。
首先，有哪些新指標？有兩種方式可以查看：
- 使用 Nsight Compute：與使用 nvprof 一樣，您可以查詢可用的指標。新工具為開發者提供了更多的指標。這些指標將會針對您正在使用的 GPU 設備顯示。在 Nsight Compute 中，同系列的設備所可用的指標都應該相同。如果您有多個不同的 GPU ，您可以選擇您希望顯示的設備。您也可以將輸出結果保存為文件。
nv-nsight-cu-cli --devices 0 --query-metrics >my_metrics.txt
(如果您需要指定完整路徑，見下文)。您也可以在命令中查詢任何架構的指標，而不受限於你使用了什么設備。
- 查看文檔：Nsight Compute 文檔在這里。Nsight Compute 文檔的另一個入口在 CUDA 文檔的 Tool 部分，你可以在側邊欄中找到。Nsight Compute CLI 文檔中常用的一個部分是 Nvprof Transition Guide（從 CUDA 10.1 Update 2 和 Nsight Compute 2019.4 開始，現在還提供 Visual Profiler 轉換指南）。該指南中有一個指標的對照表，您可以快速找到 nvvp 或 nvprof 所對應的新指標。但其中顯然沒有
gld_efficiency和gst_efficiency，所以我們需要其他方法。
當為了滿足該代碼的訪問請求，所發生的顯存(或緩存)傳輸數量最少時，則可以使可以使全局讀寫的性能最大化。對於每個線程 32-bit 數量的全局讀取請求，例如示例代碼中從 A 和 B 進行的讀取操作，我們需要 128B 來滿足每個 warp 范圍的每次請求。因此，當我們知道每個請求的最佳傳輸傳輸是多少時，監測每個請求的傳輸數量就可以得到與 gld_efficiency 和 gst_efficiency metrics 類似的效果。對於 Maxwell 及更新的 GPU ，通常滿足某 warp 一次 128-byte 的請求，最少需要 4 條傳輸（每個 32-byte）。如果數量高於此，說明並沒達到最佳性能。
不幸的是，在新工具中也沒有與以前 gld_transactions_per_request 或 gst_transactions_per_request 對應的新指標。但這些指數本質上是由分子為傳輸總數，分母為請求總數所組成的分數。至少對於計算能力7.0及以上的架構（目前為 Volta 和 Turing ），我們可以找到可以用於表示分子和分母的指標（見上述過渡指南中的對照表）。對於全局讀取傳輸，我們使用 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum ，而對於全局讀取請求則使用 l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 。如果你想知道指標的命名規則，可以在此文檔中查看。簡而言之，句點之前名稱表示了該指標所顯示的是哪個數據的信息，句點之后則顯示了數字的統計方式。對於 Volta 和更新版本的大多數指標，基本名稱和后綴（如果 .sum、.avg ...）共同組成了指標的實際名稱。一旦您了解此概念，您可以輕松的將其應用在此架構上幾乎其他任何指標上。
為什么要更改指標？Nsight Compute 的設計理念是更詳細地展示每個 GPU 的架構和顯存系統。提供了更多性能指標，更詳細地映射特定架構的特征。可自定義的 analysis section and rules 還提供了一種靈活的機制來結合多種分析數據，以構建更高級的 analyzer 。
下圖顯示了一個帶有各種指標的 GPU 顯存模型：
-
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,.per_second, l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum -
l1tex__t_bytes_pipe_lsu_mem_global_op_st.sum, .per_second -
l1tex__t_sectors_pipe_lsu_mem_local_op_ld.sum, .per_second -
l1tex__t_sectors_pipe_lsu_mem_local_op_st.sum, .per_second -
smsp__inst_executed_op_shared_ld.sum, .per_second -
smsp__inst_executed_op_shared_st.sum, .per_second -
lts__t_sectors_srcunit_tex_op_read.sum, .per_second -
lts__t_sectors_srcunit_tex_op_write.sum, .per_second -
lts__t_sectors_aperture_sysmem_op_read.sum, .per_second -
lts__t_sectors_aperture_sysmem_op_write.sum, .per_second -
dram__bytes_read.sum, .per_second -
dram__bytes_write.sum, .per_second
上表中，每行對應了圖中的一個編號。每行中的第一個指標表示其累計的總量，將 .per_second 附加到指標上，可以將其轉換成吞吐量。例如：dram__bytes_write.sum 是累積指標， dram__bytes_write.sum.per_second 則是吞吐量指標。此表僅提供一些代表性的示例，不一定適用於每條路徑。
熟悉 Nsight Compute CLI
如果你習慣使用 nvprof 或者打算創建自動化腳本，那么 Nsight Compute CLI（命令行界面）會是最好的選擇。在此，我們會使用 Linux 環境，windows 命令行的使用應該也是類似的（安裝路徑和路徑相關設置會有所不同）。Nsight Compute tool 將會與 CUDA toolkit 10.0 及以上版本一起安裝，也可以在 https://developer.nvidia.com/nsight-compute 上下載獨立的安裝程序直接安裝 Nsight Compute tool 。在運行時，您要么將 Nsight Compute 的二進制文件路徑添加到 PATH 環境變量中，要么在執行時指定完成路徑。在 CUDA 10.1 上，完整路徑為：/usr/local/cuda/NsightCompute-2019.3/，因此如果要指定CLI的完整路徑，則使用：/usr/local/cuda/NsightCompute-2019.3/nv-nsight-cu-cli 。此時就可以運行上文所介紹的查詢指標命令。對於本博客中，我們將假設您已經將路徑添加到 PATH 變量中。
雖然這不是本博客的重點，但 Nsight Compute 提供了許多功能。首先我們可以在可執行文件上以“詳細信息模式”運行它。使用 nvcc -arch=sm_70 example.cu -o example 來編譯以上代碼，修改 -arch 以適配你的 GPU 。在此示例中使用 Volta 設備 (sm_70)，但在 Turing 設備上應該同樣可以順利運行。您無法在在較早的 GPU（例如 Kepler、Maxwell、Pascal）架構上完全遵循此示例，因為計算能力 6.x 的 GPU ，與計算能力為 7.0 及更高的 GPU 可用指標有所不同。此外，計算能力 6.0 及以下的設備不支持使用 Nsight Compute。要顯示詳細信息頁面，請嘗試以下操作：
點擊查看代碼
$ /usr/local/cuda/NsightCompute-2019.3/nv-nsight-cu-cli ./example
==PROF== Connected to process 30244
==PROF== Profiling "matrix_add_2D" - 1: 0%....50%....100% - 48 passes
Success!
==PROF== Disconnected from process 30244
[30244] example@127.0.0.1
matrix_add_2D, 2019-Jun-06 23:12:59, Context 1, Stream 7
Section: GPU Speed Of Light
----------------------------------------- --------------- ------------------------------
Memory Frequency cycle/usecond 866.22
SOL FB % 21.46
Elapsed Cycles cycle 73,170
SM Frequency cycle/nsecond 1.21
Memory [%] % 56.20
Duration usecond 60.16
SOL L2 % 53.58
SOL TEX % 60.21
SM Active Cycles cycle 68,202.96
SM [%] % 8.97
----------------------------------------- --------------- ------------------------------
Section: Compute Workload Analysis
----------------------------------------- --------------- ------------------------------
Executed Ipc Active inst/cycle 0.18
Executed Ipc Elapsed inst/cycle 0.17
Issue Slots Max % 5.00
Issued Ipc Active inst/cycle 0.18
Issue Slots Busy % 4.57
SM Busy % 9.61
----------------------------------------- --------------- ------------------------------
Section: Memory Workload Analysis
----------------------------------------- --------------- ------------------------------
Memory Throughput Gbyte/second 251.25
Mem Busy % 56.20
Max Bandwidth % 53.58
L2 Hit Rate % 89.99
Mem Pipes Busy % 3.36
L1 Hit Rate % 90.62
----------------------------------------- --------------- ------------------------------
Section: Scheduler Statistics
----------------------------------------- --------------- ------------------------------
Active Warps Per Scheduler warp 11.87
Eligible Warps Per Scheduler warp 0.15
No Eligible % 95.39
Instructions Per Active Issue Slot inst/cycle 1
Issued Warp Per Scheduler 0.05
One or More Eligible % 4.61
----------------------------------------- --------------- ------------------------------
Section: Warp State Statistics
----------------------------------------- --------------- ------------------------------
Avg. Not Predicated Off Threads Per Warp 29.87
Avg. Active Threads Per Warp 32
Warp Cycles Per Executed Instruction cycle 261.28
Warp Cycles Per Issued Instruction 257.51
Warp Cycles Per Issue Active 257.51
----------------------------------------- --------------- ------------------------------
Section: Instruction Statistics
----------------------------------------- --------------- ------------------------------
Avg. Executed Instructions Per Scheduler inst 3,072
Executed Instructions inst 983,040
Avg. Issued Instructions Per Scheduler inst 3,116.96
Issued Instructions inst 997,428
----------------------------------------- --------------- ------------------------------
Section: Launch Statistics
----------------------------------------- --------------- ------------------------------
Block Size 1,024
Grid Size 1,024
Registers Per Thread register/thread 16
Shared Memory Configuration Size byte 0
Dynamic Shared Memory Per Block byte/block 0
Static Shared Memory Per Block byte/block 0
Threads thread 1,048,576
Waves Per SM 6.40
----------------------------------------- --------------- ------------------------------
Section: Occupancy
----------------------------------------- --------------- ------------------------------
Block Limit SM block 32
Block Limit Registers block 4
Block Limit Shared Mem block inf
Block Limit Warps block 2
Achieved Active Warps Per SM warp 48.50
Achieved Occupancy % 75.78
Theoretical Active Warps per SM warp/cycle 64
Theoretical Occupancy % 100
----------------------------------------- --------------- ------------------------------
它輸出了很多內容（如果您的代碼調用了多個 kernel ，它會收集並顯示每個的狀態）。我們不會介紹所有細節，但是有些 SOL（使用率） 方面數據需要注意：計算分析、顯存分析、調度器、warp 狀態、指令、配置信息、occupancy 分析。您可以使用命令行參數來決定要顯示哪些部分。命令行參數幫助可以使用 --help 來獲得，也可以在文檔中查看。請注意，選擇輸出的部分和指標的選擇會影響分析的時間和輸出的大小。
我們可以使用以上數據來對我們的目標（全局讀取/寫入效率）進行判斷。我們以類似 nvprof 的方式獲取信息：
$ nv-nsight-cu-cli --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./example
==PROF== Connected to process 30749
==PROF== Profiling "matrix_add_2D" - 1: 0%....50%....100% - 4 passes
Success!
==PROF== Disconnected from process 30749
[30749] example@127.0.0.1
matrix_add_2D, 2019-Jun-06 23:25:45, Context 1, Stream 7
Section: Command line profiler metrics
------------------------------------------------ ------------ ------------------------------
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum request 65,536
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 2,097,152
------------------------------------------------ ------------ ------------------------------
上面第一個指標表示【請求總數】 作為分母， 第二個指標表示【傳輸總數】 作為分子。將他們相除，即可得出每個請求需要 32 個傳輸的結果。這意味着，warp 中每個線程都在進行單獨的傳輸，並沒有合並。
使用 Nsight Compute GUI
如果我們想用 GUI 來顯示這些數據怎么辦？在 linux 上，我們需要使用 X session 來啟動 GUI 程式應用。因此你需要在支持圖形功能的設備上運行。要運行 Nsight Compute GUI ，在命令行輸入以下命令：
/usr/local/cuda/NsightCompute-2019.3/nv-nsight-cu
如果你已經將路徑加入到 PATH 變量中，可以直接輸入 nv-nsight-cu ，接下來你就會看到 Nsight Compute 的圖形界面：
從最簡單開始，點擊 Quick Launch 下的 Continue （或者你可以點擊 New Project 下的 Create New Project 來創建一個新的項目。）接下來，應該會打開一個分析配置窗口，點擊窗口下方的 Additional Options ，點擊 Other 分頁，接下來填入 Application Executable: 、 Output File: 、 Metrics: ：
在此我們輸入了要分析的可執行文件的路徑和名稱（示例）、導出分析結果的文件和用逗號分隔開的指標：
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum
之后，您可以最小化 Additional Options 窗口，並點擊藍色的 Launch 按鈕。然后分析器就會運行並捕獲數據，結果如下：
在上圖中，顯示了所填指標的數據，及其他資料的數據（在本例中為顯存負載）。請注意，在此狀態下保存到磁盤的文件是人類不可讀的，需要在 Nsight Compute GUI 中查看。對於人類可讀文件，在大多報告頁面中都有導出按鈕，通常位於右上角。
如果您想更詳細地探索 GUI 功能，文檔包含介紹 GUI 的快速入門部分。
修復代碼
示例代碼的執行效率低（每個請求的傳輸數量過多）的原因是由於我們使用了 2D 索引：
... C[idx][idy] = A[idx][idy] + B[idx][idy];
使用 threadIdx.x（即 idx）構建的索引應該出現在最后一個下標中，以便跨 warp 進行合並訪問；相反，它出現在第一個下標中。雖然兩種方法都可以給出正確的結果，但他們的性能差距很大。這種設計會倒是 warp 中每個線程訪問的是顯存中 “column” 的數據，而不是 “row”（即相鄰）的數據。我們可以修改 kernel 代碼來優化這個問題：
__global__ void matrix_add_2D(const arr_t * __restrict__ A, const arr_t * __restrict__ B, arr_t * __restrict__ C, const size_t sw, const size_t sh){
size_t idx = threadIdx.x+blockDim.x*(size_t)blockIdx.x;
size_t idy = threadIdx.y+blockDim.y*(size_t)blockIdx.y;
if ((idy < sh) && (idx < sw)) C[idy][idx] = A[idy][idx] + B[idy][idx];
}
唯一改變的只有最后一行代碼，我們交換了 idx 和 idy 的位置。當我們重新編譯修改后的代碼並運行上面相同的分析時，可以看到：
$ nv-nsight-cu-cli --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./example
==PROF== Connected to process 5779
==PROF== Profiling "matrix_add_2D" - 1: 0%....50%....100% - 4 passes
Success!
==PROF== Disconnected from process 5779
[5779] example@127.0.0.1
matrix_add_2D, 2019-Jun-11 12:01:26, Context 1, Stream 7
Section: Command line profiler metrics
----------------------------------------------- --------------- ------------
l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum request 65,536
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 262,144
----------------------------------------------- --------------- ------------
現在這個比例變為了 4:1 （每個請求的傳輸數量），說明達到了 32 byte 的期望值，並且讀取（和寫入）的效率也比以前高很多。、
由於這項工作涉及將新舊結果的比較，因此我們繼續演示 GUI 的附加功能。我們可以使用 GUI 收集兩種情況的分析結果，並進行比較。如上文所述，我們先收集舊代碼的數據，並保持 GUI 界面打開 ，然后選擇左上角的 Connect 按鈕，將輸出文件改為新名稱，並把分析文件改成修改后的文件。然后點擊 Launch 創建一個新選項卡，其中的數據則是修改后的代碼的。最后回到初始結果選項卡，選擇頂部的 Add Baseline 並選擇新結果的選項卡，即可查看各項指標的差異。
在本例中，我們看到改進后的傳輸次數相比原始的減少到了1/8
提高了顯存的使用效率，將會提高使用此顯存的代碼的性能，這意味着程序運行的速度也會更快。我們可以使用上一篇博客總介紹的 Nsight Systems profiler ，檢查更改前后內核的運行時間來驗證這一點。我們運行以下命令，這與我們上一篇博客中減少的第一個 CLI 命令類似：
$ nsys profile -o example.nsysprofout --stats=true ./example
但是本文的重點是 Nsight Compute ，我們可以通過 GPU SOL 報告中的 Elapsed Cycles 部分來進行類似的檢測。我們繼續使用上一節介紹的比較功能。在 GUI 中，首先選擇左上角的 Connect 按鈕打開配置設置，選擇 Additional Options 下拉菜單，你可以將 Other 中的指標都清楚，然后選擇 Sections 選項卡，選擇 GPU Speed of Light（可以將其他選項取消，以簡化輸出並減少分析時間）。您可能還需要更改分析輸出文件的文件名。然后點擊藍色的 launch 按鈕。
單擊 啟動 按鈕以收集新的分析數據。與先前一樣，我們對原版和修改版重復執行這些步驟。然后將原版設為基准，並查看對比。
如上圖所述，我們可以看到修改后的代碼的執行時間減少了約 68%。包括其他各項數據，也顯示此更改對性能的影響。
有哪些新功能？
與 NVIDIA Visual Profiler 和 nvprof 相比，Nsight Compute 中有許多新功能，我們在本博客中僅涉及其中的一些。
Nsight Compute GUI 與 Visual Profiler 相比的新功能：
- 在工具中比較分析結果
- 交互式分析模式（使用 API 流和參數捕獲）
- 具有跨操作系統支持的遠程操作
Nsight Compute GUI 和 CLI 與 Visual Profiler/nvprof 相比的新功能：
- 更詳細的指標
- 可自定義指標和基於 python 的分析引導
- 更穩定的數據收集（時鍾控制、緩存重置……）
- 減少內核重運行的開銷（與第一次的差異）
- 支持新的 CUDA/NVTX 功能（例如圖形支持、nvtx 過濾器描述）
結論
與 nvprof 和 Visual Profiler 相比，新工具旨在提供相同（並且更好）的功能，但需要一些新設置和新方法才能獲得相似的結果。關於作為本博客主要關注點的指標分析，熟悉新指標很重要，將新指標進行組合來獲取您想要的結果。