$ NSYS stats --help-reports
The following built-in reports are available:
cuda_api_gpu_sum[:nvtx-name][:base|:mangled] -- CUDA Summary (API/Kernels/MemOps)
cuda_gpu_sum[:nvtx-name][:base|:mangled] -- CUDA GPU Summary (Kernels/MemOps)
cuda_gpu_trace[:nvtx-name][:base|:mangled] -- CUDA GPU Trace，带 device duration 非常有用
cuda_api_sum -- CUDA API Summary
cuda_api_trace -- CUDA API Trace
cuda_gpu_kern_gb_sum[:nvtx-name][:base|:mangled] -- CUDA GPU Kernel/Grid/Block Summary
cuda_gpu_kern_sum[:nvtx-name][:base|:mangled] -- CUDA GPU Kernel Summary
cuda_gpu_mem_size_sum -- CUDA GPU MemOps Summary (by Size)
cuda_gpu_mem_time_sum -- CUDA GPU MemOps Summary (by Time)
cuda_green_context_sum -- CUDA Green Context Summary
cuda_kern_exec_sum[:nvtx-name][:base|:mangled] -- CUDA Kernel Launch & Exec Time Summary，带 DevId
cuda_kern_exec_trace[:nvtx-name][:base|:mangled] -- CUDA Kernel Launch & Exec Time Trace，带 DevId
dx11_pix_sum -- DX11 PIX Range Summary
dx12_gpu_marker_sum -- DX12 GPU Command List PIX Ranges Summary
dx12_pix_sum -- DX12 PIX Range Summary
mpi_event_sum -- MPI Event Summary
mpi_event_trace -- MPI Event Trace
mpi_msg_size_sum -- MPI Message Size Summary
network_congestion[:ticks_threshold=<ticks_per_ms>] -- Network Devices Congestion
nvtx_gpu_proj_sum -- NVTX GPU Projection Summary
nvtx_gpu_proj_trace -- NVTX GPU Projection Trace
nvtx_kern_sum[:base|:mangled] -- NVTX Range Kernel Summary
nvtx_pushpop_sum -- NVTX Push/Pop Range Summary
nvtx_pushpop_trace -- NVTX Push/Pop Range Trace
nvtx_startend_sum -- NVTX Start/End Range Summary
nvtx_sum -- NVTX Range Summary
nvvideo_api_sum -- NvVideo API Summary
openacc_sum -- OpenACC Summary
opengl_khr_gpu_range_sum -- OpenGL KHR_debug GPU Range Summary
opengl_khr_range_sum -- OpenGL KHR_debug Range Summary
openmp_sum -- OpenMP Summary
osrt_sum -- OS Runtime Summary
syscall_sum -- Syscall Summary
um_cpu_page_faults_sum -- Unified Memory CPU Page Faults Summary
um_sum[:rows=<limit>] -- Unified Memory Analysis Summary
um_total_sum -- Unified Memory Totals Summary
vulkan_api_sum -- Vulkan API Summary
vulkan_api_trace -- Vulkan API Trace
vulkan_gpu_marker_sum -- Vulkan GPU Range Summary
vulkan_marker_sum -- Vulkan Range Summary
wddm_queue_sum -- WDDM Queue Utilization Summary
For more information, use '--help-reports <report_name>'
Time : Percentage of "Total Time"
Total Time : Total time used by all executions of this kernel
Instances : Number of executions of this kernel
Avg : Average execution time of this kernel
Med : Median execution time of this kernel
Min : Smallest execution time of this kernel
Max : Largest execution time of this kernel
StdDev : Standard deviation of execution time of this kernel
**Category** : Category of the operation
**Operation** : Name of the kernel
Time Total Time Instances Avg Med Min Max StdDev Category Operation
...
0.2% 38.495 ms 64 601.484 μs 600.532 μs 580.485 μs 623.235 μs 15.708 μs CUDA_KERNEL ampere_fp16_s16816gemm_fp16_128x128_ldg8_f2f_stages_32x5_tn
0.2% 36.876 ms 1 36.876 ms 36.876 ms 36.876 ms 36.876 ms 0 ns CUDA_API cudaProfilerStop
0.1% 30.602 ms 32768 933 ns 832 ns 701 ns 495.466 μs 4.397 μs CUDA_API cudaEventDestroy
0.1% 28.753 ms 33536 857 ns 861 ns 561 ns 34.715 μs 452 ns CUDA_API cuKernelGetFunction
0.1% 24.788 ms 49088 505 ns 480 ns 448 ns 2.016 μs 160 ns MEMORY_OPER [CUDA memset]
...
...
Time Total Time Instances Avg Med Min Max StdDev Category Operation
0.3% 45.617 ms 33280 1.370 μs 1.376 μs 1.312 μs 1.952 μs 33 ns CUDA_KERNEL void at::native::vectorized_elementwise_kernel<(int)4, at::native::rsqrt_kernel_cuda(at::TensorIteratorBase &)::[lambda() (instance 2)]::operator ()() const::[lambda() (instance 2)]::operator ()() const::[lambda(float) (instance 1)], std::array<char *, (unsigned long)2>>(int, T2, T3)
0.3% 43.317 ms 33280 1.301 μs 1.312 μs 1.248 μs 2.016 μs 40 ns CUDA_KERNEL void at::native::vectorized_elementwise_kernel<(int)4, at::native::CUDAFunctorOnSelf_add<float>, std::array<char *, (unsigned long)2>>(int, T2, T3)
0.2% 38.495 ms 64 601.484 μs 600.532 μs 580.485 μs 623.235 μs 15.708 μs CUDA_KERNEL ampere_fp16_s16816gemm_fp16_128x128_ldg8_f2f_stages_32x5_tn
0.1% 24.788 ms 49088 505 ns 480 ns 448 ns 2.016 μs 160 ns MEMORY_OPER [CUDA memset]
0.1% 13.702 ms 64 214.094 μs 214.161 μs 206.562 μs 221.122 μs 2.928 μs CUDA_KERNEL void pytorch_flash::flash_fwd_kernel<pytorch_flash::Flash_fwd_kernel_traits<(int)128, (int)64, (int)64, (int)4, (bool)0, (bool)0, cutlass::half_t, pytorch_flash::Flash_kernel_traits<(int)128, (int)64, (int)64, (int)4, cutlass::half_t>>, (bool)0, (bool)1, (bool)0, (bool)0, (bool)1, (bool)1, (bool)0>(pytorch_flash::Flash_fwd_params)
...
Output: All time values default to nanoseconds
Start : Timestamp of start time
Duration : Length of event
CorrId : Correlation ID
GrdX, GrdY, GrdZ : Grid values
BlkX, BlkY, BlkZ : Block values
Reg/Trd : Registers per thread
StcSMem : Size of Static Shared Memory
DymSMem : Size of Dynamic Shared Memory
Bytes : Size of memory operation
Throughput : Memory throughput
SrcMemKd : Memcpy source memory kind or memset memory kind
DstMemKd : Memcpy destination memory kind
Device : GPU device name and ID
Ctx : Context ID
GreenCtx: Green context ID
Strm : Stream ID
Name : Trace event name
...
Start Duration CorrId GrdX GrdY GrdZ BlkX BlkY BlkZ Reg/Trd StcSMem DymSMem Bytes Throughput SrcMemKd DstMemKd Device Ctx GreenCtx Strm Name
0.0126451s 2.048 μs 728906 - - - - - - - - - 1 B 476.84 KiB/s Device Pinned NVIDIA A40 (0) 1 - 7 [CUDA memcpy Device-to-Host]
0.0128317s 1.952 μs 728906 - - - - - - - - - 1 B 500.29 KiB/s Device Pinned NVIDIA A40 (1) 1 - 7 [CUDA memcpy Device-to-Host]
0.0127674s 5.121 μs 728917 2 1 1 128 4 1 36 0 B 8.00 KiB - - - - NVIDIA A40 (0) 1 - 7 void at::native::tensor_kernel_scan_innermost_dim<long, std::plus<long>>(T1 *, const T1 *, unsigned int, unsigned int, unsigned int, T1, T2)
0.0129518s 5.088 μs 728917 2 1 1 128 4 1 36 0 B 8.00 KiB - - - - NVIDIA A40 (1) 1 - 7 void at::native::tensor_kernel_scan_innermost_dim<long, std::plus<long>>(T1 *, const T1 *, unsigned int, unsigned int, unsigned int, T1, T2)
0.0128081s 1.632 μs 728927 4 1 1 128 1 1 32 0 B 0 B - - - - NVIDIA A40 (0) 1 - 7 void at::native::vectorized_elementwise_kernel<(int)2, at::native::CUDAFunctorOnSelf_add<long>, std::array<char *, (unsigned long)2>>(int, T2, T3)
0.012992s 1.600 μs 728927 4 1 1 128 1 1 32 0 B 0 B - - - - NVIDIA A40 (1) 1 - 7 void at::native::vectorized_elementwise_kernel<(int)2, at::native::CUDAFunctorOnSelf_add<long>, std::array<char *, (unsigned long)2>>(int, T2, T3)
...
Time Total Time Num Calls Avg Med Min Max StdDev Name
60.5% 4.217 s 660690 6.382 μs 5.380 μs 3.416 μs 3.777 ms 16.781 μs cudaLaunchKernel
11.4% 794.180 ms 1032 769.553 μs 7.840 μs 2.274 μs 375.076 ms 16.501 ms cudaStreamSynchronize
4.7% 326.003 ms 33536 9.721 μs 8.276 μs 6.432 μs 2.721 ms 51.003 μs cuLaunchKernel
4.6% 320.383 ms 170982 1.873 μs 1.453 μs 390 ns 471.301 μs 2.189 μs cudaEventQuery
4.4% 308.913 ms 49088 6.293 μs 5.801 μs 5.059 μs 767.245 μs 4.862 μs cudaMemsetAsync
4.3% 301.013 ms 32768 9.186 μs 8.917 μs 8.085 μs 126.006 μs 1.914 μs cuLaunchKernelEx
3.8% 264.435 ms 196608 1.345 μs 1.162 μs 490 ns 363.048 μs 1.606 μs cudaEventRecord
2.6% 182.189 ms 131074 1.390 μs 1.422 μs 611 ns 737.258 μs 2.206 μs cudaStreamWaitEvent
0.9% 59.757 ms 49664 1.203 μs 1.172 μs 802 ns 37.179 μs 660 ns cudaStreamIsCapturing
0.7% 48.591 ms 32768 1.482 μs 1.343 μs 1.142 μs 1.204 ms 7.098 μs cudaEventCreateWithFlags
0.5% 36.876 ms 1 36.876 ms 36.876 ms 36.876 ms 36.876 ms 0 ns cudaProfilerStop
0.4% 30.602 ms 32768 933 ns 832 ns 701 ns 495.466 μs 4.397 μs cudaEventDestroy
0.4% 28.753 ms 33536 857 ns 861 ns 561 ns 34.715 μs 452 ns cuKernelGetFunction
0.3% 24.062 ms 32768 734 ns 721 ns 631 ns 31.939 μs 355 ns cudaStreamGetCaptureInfo_v2
0.2% 16.194 ms 1544 10.488 μs 9.948 μs 6.783 μs 46.407 μs 3.259 μs cudaMemcpyAsync
0.1% 9.890 ms 32768 301 ns 290 ns 250 ns 32.631 μs 425 ns cudaGetFuncBySymbol
0.0% 247.242 μs 2 123.621 μs 123.621 μs 121.838 μs 125.404 μs 2.521 μs cudaDeviceSynchronize
0.0% 23.403 μs 2 11.701 μs 11.701 μs 10.499 μs 12.904 μs 1.700 μs cuProfilerStart
0.0% 2.164 μs 2 1.082 μs 1.082 μs 320 ns 1.844 μs 1.077 μs cudaThreadExchangeStreamCaptureMode
Start Duration Name Result CorrID Pid Tid T-Pri Thread Name
0.010916s 12.904 μs cuProfilerStart 0 728849 1261445 1261445 20 python
0.0109509s 10.499 μs cuProfilerStart 0 728849 1261444 1261444 20 python
0.0122638s 11.722 μs cudaEventQuery 0 728853 1261444 1261444 20 python
0.0123082s 34.575 μs cudaMemcpyAsync 0 728860 1261444 1261444 20 python
0.0123437s 53.671 μs cudaStreamSynchronize 0 728861 1261444 1261444 20 python
...
Time Total Time Instances Avg Med Min Max StdDev GridXYZ BlockXYZ Name
18.7% 3.091 s 32640 94.709 μs 21.728 μs 12.704 μs 2.066 ms 130.244 μs 2 1 1 512 1 1 ncclDevKernel_AllReduce_Sum_f16_RING_LL(ncclDevKernelArgsStorage<(unsigned long)4096>)
17.0% 2.810 s 32640 86.100 μs 86.080 μs 65.217 μs 108.609 μs 10.369 μs 168 2 1 512 1 1 void at::native::<unnamed>::CatArrayBatchedCopy_contig<at::native::<unnamed>::OpaqueType<(unsigned int)2>, unsigned int, (int)4, (int)128, (int)1>(T1 *, at::native::<unnamed>::CatArrInputTensorMetadata<T1, T2, T4, T5>, at::native::<unnamed>::TensorSizeStride<T2, (unsigned int)4>, int, T2)
16.2% 2.675 s 32640 81.943 μs 81.952 μs 80.769 μs 84.096 μs 430 ns 8 43 1 32 1 1 void cutlass::Kernel2<cutlass_80_wmma_tensorop_f16_s161616gemm_f16_16x16_128x2_tn_align8>(T1::Params)
11.7% 1.939 s 48960 39.605 μs 39.360 μs 38.560 μs 43.841 μs 609 ns 16 1 3 128 1 1 ampere_fp16_s16816gemm_fp16_128x64_ldg8_f2f_stages_64x3_tn
8.5% 1.408 s 16320 86.252 μs 86.241 μs 84.897 μs 88.161 μs 426 ns 64 1 1 128 1 1 ampere_fp16_s16816gemm_fp16_64x64_sliced1x2_ldg8_f2f_stages_64x6_tn
Time Total Time Instances Avg Med Min Max StdDev Name
20.3% 3.362 s 32768 102.589 μs 21.952 μs 12.704 μs 2.556 ms 181.043 μs ncclDevKernel_AllReduce_Sum_f16_RING_LL(ncclDevKernelArgsStorage<(unsigned long)4096>)
17.0% 2.810 s 32640 86.100 μs 86.080 μs 65.217 μs 108.609 μs 10.369 μs void at::native::<unnamed>::CatArrayBatchedCopy_contig<at::native::<unnamed>::OpaqueType<(unsigned int)2>, unsigned int, (int)4, (int)128, (int)1>(T1 *, at::native::<unnamed>::CatArrInputTensorMetadata<T1, T2, T4, T5>, at::native::<unnamed>::TensorSizeStride<T2, (unsigned int)4>, int, T2)
16.2% 2.675 s 32640 81.943 μs 81.952 μs 80.769 μs 84.096 μs 430 ns void cutlass::Kernel2<cutlass_80_wmma_tensorop_f16_s161616gemm_f16_16x16_128x2_tn_align8>(T1::Params)
11.7% 1.939 s 48960 39.605 μs 39.360 μs 38.560 μs 43.841 μs 609 ns ampere_fp16_s16816gemm_fp16_128x64_ldg8_f2f_stages_64x3_tn
8.5% 1.408 s 16320 86.252 μs 86.241 μs 84.897 μs 88.161 μs 426 ns ampere_fp16_s16816gemm_fp16_64x64_sliced1x2_ldg8_f2f_stages_64x6_tn
Output: All time values default to nanoseconds
PID : Process ID that made kernel launch call
TID : Thread ID that made kernel launch call
DevId : CUDA Device ID that executed kernel (which GPU)
Count : Number of kernel records
QCount : Number of kernel records with positive queue time
Average, Median, Minimum, Maximum, and Standard Deviation for:
TAvg, TMed, TMin, TMax, TStdDev : Total time
AAvg, AMed, AMin, AMax, AStdDev : API time
QAvg, QMed, QMin, QMax, QStdDev : Queue time
KAvg, KMed, KMin, KMax, KStdDev : Kernel time
API Name : Name of CUDA API call used to launch kernel
Kernel Name : Name of CUDA Kernel
PID TID DevId Count QCount TAvg TMed TMin TMax TStdDev AAvg AMed AMin AMax AStdDev QAvg QMed QMin QMax QStdDev KAvg KMed KMin KMax KStdDev API Name Kernel Name
1261444 1261444 0 192 192 217.507 ms 222.078 ms 1.441 ms 385.721 ms 122.846 ms 113.625 μs 8.440 μs 6.913 μs 2.721 ms 468.758 μs 216.289 ms 220.967 ms 772.543 μs 384.482 ms 122.694 ms 1.104 ms 1.101 ms 653.861 μs 1.570 ms 446.898 μs cuLaunchKernel void cutlass::Kernel2<cutlass_80_tensorop_f16_s16816gemm_relu_f16_256x128_32x3_tn_align8>(T1::Params)
1261445 1261445 1 192 192 217.177 ms 221.690 ms 1.452 ms 385.723 ms 122.826 ms 112.853 μs 8.175 μs 6.432 μs 2.652 ms 461.644 μs 215.908 ms 220.519 ms 756.472 μs 384.433 ms 122.673 ms 1.156 ms 1.149 ms 679.683 μs 1.645 ms 470.911 μs cuLaunchKernel void cutlass::Kernel2<cutlass_80_tensorop_f16_s16816gemm_relu_f16_256x128_32x3_tn_align8>(T1::Params)
1261444 1261444 0 32 32 215.004 ms 219.358 ms 3.718 ms 379.305 ms 123.986 ms 94.487 μs 8.826 μs 8.355 μs 590.484 μs 202.555 μs 214.323 ms 218.764 ms 3.123 ms 378.171 ms 123.872 ms 586.407 μs 585.669 μs 580.485 μs 593.093 μs 3.210 μs cudaLaunchKernel ampere_fp16_s16816gemm_fp16_128x128_ldg8_f2f_stages_32x5_tn
1261445 1261445 1 32 32 214.674 ms 218.946 ms 3.810 ms 379.351 ms 123.947 ms 84.002 μs 8.570 μs 8.115 μs 612.235 μs 201.565 μs 213.974 ms 218.318 ms 3.189 ms 378.134 ms 123.845 ms 616.561 μs 617.523 μs 607.971 μs 623.235 μs 4.673 μs cudaLaunchKernel ampere_fp16_s16816gemm_fp16_128x128_ldg8_f2f_stages_32x5_tn
1261444 1261444 0 32 32 214.570 ms 218.851 ms 3.243 ms 379.380 ms 124.088 ms 82.630 μs 7.694 μs 7.164 μs 600.192 μs 197.347 μs 214.273 ms 218.627 ms 3.017 ms 378.573 ms 123.987 ms 214.890 μs 214.353 μs 206.562 μs 221.122 μs 3.088 μs cudaLaunchKernel void pytorch_flash::flash_fwd_kernel<pytorch_flash::Flash_fwd_kernel_traits<(int)128, (int)64, (int)64, (int)4, (bool)0, (bool)0, cutlass::half_t, pytorch_flash::Flash_kernel_traits<(int)128, (int)64, (int)64, (int)4, cutlass::half_t>>, (bool)0, (bool)1, (bool)0, (bool)0, (bool)1, (bool)1, (bool)0>(pytorch_flash::Flash_fwd_params)
API Start API Dur Queue Start Queue Dur Kernel Start Kernel Dur Total Dur PID TID DevId API Function GridXYZ BlockXYZ Kernel Name
0.0125269s 38.221 μs - - 0.0125613s 2.016 μs 38.221 μs 1261444 1261444 0 cudaLaunchKernel 1 1 1 128 1 1 void at::native::elementwise_kernel<(int)128, (int)4, void at::native::gpu_kernel_impl_nocast<at::native::AUnaryFunctor<long, long, bool, at::native::<unnamed>::CompareEqFunctor<long>>>(at::TensorIteratorBase &, const T1 &)::[lambda(int) (instance 1)]>(int, T3)
0.0126038s 14.777 μs - - 0.0126172s 3.200 μs 16.579 μs 1261444 1261444 0 cudaLaunchKernel 1 1 1 8 1 1 void at::native::reduce_kernel<(int)512, (int)1, at::native::ReduceOp<bool, at::native::func_wrapper_t<bool, at::native::or_kernel_cuda(at::TensorIterator &)::[lambda() (instance 1)]::operator ()() const::[lambda() (instance 12)]::operator ()() const::[lambda(bool, bool) (instance 1)]>, unsigned int, bool, (int)4, (int)4>>(T3)
0.0127125s 40.926 μs - - 0.0127485s 2.080 μs 40.926 μs 1261445 1261445 1 cudaLaunchKernel 1 1 1 128 1 1 void at::native::elementwise_kernel<(int)128, (int)4, void at::native::gpu_kernel_impl_nocast<at::native::AUnaryFunctor<long, long, bool, at::native::<unnamed>::CompareEqFunctor<long>>>(at::TensorIteratorBase &, const T1 &)::[lambda(int) (instance 1)]>(int, T3)
0.0127501s 18.144 μs - - 0.0127674s 5.121 μs 22.386 μs 1261444 1261444 0 cudaLaunchKernel 2 1 1 128 4 1 void at::native::tensor_kernel_scan_innermost_dim<long, std::plus<long>>(T1 *, const T1 *, unsigned int, unsigned int, unsigned int, T1, T2)
0.0127882s 14.357 μs - - 0.0128009s 3.136 μs 15.884 μs 1261445 1261445 1 cudaLaunchKernel 1 1 1 8 1 1 void at::native::reduce_kernel<(int)512, (int)1, at::native::ReduceOp<bool, at::native::func_wrapper_t<bool, at::native::or_kernel_cuda(at::TensorIterator &)::[lambda() (instance 1)]::operator ()() const::[lambda() (instance 12)]::operator ()() const::[lambda(bool, bool) (instance 1)]>, unsigned int, bool, (int)4, (int)4>>(T3)
-
NVTX_EVENTS
start end eventType rangeId category color text **globalTid** endGlobalTid textId domainId uint64Value int64Value doubleValue uint32Value int32Value floatValue jsonTextId jsonText binaryData -
CUPTI_ACTIVITY_KIND_KERNEL
没有 globalTid，因为 KERNEL 是 GPU 端的事件，本质上不知道哪个 CPU 线程把它扔进队列的。
GPU 不关心 launch 来自哪个线程，它只看到：
- 这个 kernel 在哪个 context（context 是按进程创建的）
- 在哪个 stream 上
- 在哪个 device 上
所以 KERNEL 表里有的是进程级和设备级的字段：deviceId — 哪张卡、contextId — 哪个 CUDA context、streamId — 哪个 stream、globalPid — 哪个进程。
而线程级的信息只在 CPU 端的 RUNTIME 表里有（globalTid），因为 cudaLaunchKernel 是某个 CPU 线程调用的 API。
start end deviceId contextId greenContextId streamId correlationId **globalPid** demangledName shortName mangledName launchType cacheConfig registersPerThread gridX gridY gridZ blockX blockY blockZ staticSharedMemory dynamicSharedMemory localMemoryPerThread localMemoryTotal gridId sharedMemoryExecuted graphNodeId sharedMemoryLimitConfig qmdBulkReleaseDone qmdPreexitDone qmdLastCtaDone graphId clusterX clusterY clusterZ clusterSchedulingPolicy maxPotentialClusterSize maxActiveClusters sharedMemoryRequestedPercentage tensorSizeMinusOneElements -
CUPTI_ACTIVITY_KIND_RUNTIME
start end eventClass **globalTid** correlationId nameId returnValue callchainId -
globalTid globalPid 具体构成
globalTid 把 (VMID, PID, TID) 三个信息打包到一个 64 位整数里。
bit 63 bit 0 [ VMID ][ PID (24b) ][ TID (24b) ]- 低 24 位 = Linux TID（LWP id，nsys GUI 里的 thread id）
- 中间 24 位 = Linux PID（nsys GUI 里的 process id）
- 高位 = VMID（一般是 1，native 环境固定值）
globalPid 的 TID 部分是 0。可以查 sqlite 数据库验证：
python -c " import sqlite3 conn = sqlite3.connect('/work/hdd/bgnm/group4/tp4_bs1_seq512.sqlite') print('RUNTIME distinct globalTids:') for r in conn.execute('SELECT DISTINCT globalTid FROM CUPTI_ACTIVITY_KIND_RUNTIME WHERE globalTid IS NOT NULL'): print(f' {r[0]:#x} ({r[0]})') print('KERNEL distinct globalPids:') for r in conn.execute('SELECT DISTINCT globalPid FROM CUPTI_ACTIVITY_KIND_KERNEL WHERE globalPid IS NOT NULL'): print(f' {r[0]:#x} ({r[0]})') " -
提取 PID TID
python -c " gid = 0x1134671134802 print(f'PID = {(gid >> 24) & 0xFFFFFF}') print(f'TID = {gid & 0xFFFFFF}') "
这是 --gpu-metrics-device=all 开启后 nsys 通过 GPU 硬件性能计数器采样得到的数据,10 kHz 表示每秒 10000 次采样(每 100μs 一次)。这些是分析 GPU 利用情况最硬核的指标。我按图里的分组讲。
在 Nsight Systems (nsys) 的 GPU Metrics 中，GPC Clock 和 SYS Clock 分别代表了 GPU 内部两个完全独立的物理时钟域 (Clock Domains)。它们分别驱动着芯片上不同职能的硬件模块。
-
GPC Clock Frequency (图形处理集群时钟)
物理对应： GPC 全称是 Graphics Processing Cluster。这是 GPU 上真正的“计算肌肉”。
覆盖的硬件： 它驱动着所有的 SM (流式多处理器)。这意味着你代码里用到的 CUDA Cores (ALU)、Tensor Cores、L1 Cache / Shared Memory、Warp 调度器、以及纹理单元，全部都是跟着 GPC Clock 的节奏在跳动。
核心职能： 纯粹的数学运算和逻辑指令执行（如 FMA、矩阵乘加等）。
动态表现： 当你的 Kernel 是计算密集型 (Compute-Bound) 的时候，GPC Clock 会迅速飙升到最高（Boost 频率），疯狂算数。当 GPU 空闲时，为了省电，它会大幅度降频。
-
SYS Clock Frequency (系统时钟)
物理对应： SYS 指的是 GPU 的 System Fabric / Interconnect 层，也就是连接各个计算单元和存储单元的全局大网。
覆盖的硬件： 它驱动着 L2 Cache (二级缓存)、Crossbar (全局交叉开关互连网)、前端命令处理器 (CP)、以及显存控制器 (Memory Controllers) 的接口侧。
核心职能： 数据搬运和全局路由。它负责把数据从显存搬进 L2 Cache，再通过 Crossbar 分发给各个 GPC 里的 SM；同时负责各个 SM 之间的协同通信。
动态表现： 只要 GPU 内部存在大量的数据流转（无论是 SM 在疯狂读写 L2，还是 Host 正在往 Device 拷贝数据），SYS Clock 就会保持在高位。
实战用途:如果你发现 kernel 跑得慢,先看 GPC Clock 是不是被降频了。常见原因:温度过高、功耗墙、nvidia-smi 锁了较低频率。A40 的 boost clock 应该接近 1.7GHz,如果只有 1.3GHz 跑那就是降频了。
这一组告诉你 GPU 不同引擎在不在工作:
| 指标 | 含义 |
|---|---|
| GPU Active | GPU 整体是否在工作(任意引擎工作即为 active)。这是最高层的"GPU 在干活"指标 |
| Sync Copy Engine Active | 同步拷贝引擎在工作(cudaMemcpy 那种) |
| Async Copy Engine Active | 异步拷贝引擎在工作(cudaMemcpyAsync)。这两个引擎可以和 compute 并行,所以 H2D/D2H 和 kernel 可以重叠 |
| GR Active | Graphics Engine active,跑 kernel 时这个会激活(对纯 compute 程序来说,GR Active ≈ kernel 在跑) |
关键概念:GPU 上有多个独立引擎(SM、copy engine 等),它们可以并行工作。所以理想情况下,H2D 拷贝(Async Copy Engine)和 kernel 计算(GR Active)可以同时发生——这就是 stream 重叠的基础。
实战用途:看 H2D/D2H 和 kernel 是不是真的在重叠。如果你的代码用了 multi-stream 但 Async Copy Engine 和 GR Active 从来不同时激活,说明重叠没生效。
| 指标 | 含义 |
|---|---|
| Sync Compute in Flight | 同步上下文里有多少 compute work 在跑 |
| Async Compute in Flight | 异步上下文里有多少 compute work 在跑 |
对你这种 PyTorch CUDA 程序,大部分活动会在 Async 这边。
| 指标 | 含义 |
|---|---|
| SMs Active | 有百分之多少的 SM 在工作。例如 A40 有 84 个 SM,如果只用了 42 个,SMs Active = 50% |
- GPU Active = GPU 上任何部件在工作(可能只是 copy engine 在拷贝)
- SMs Active = SM 计算单元在工作(真在算东西)
LLM decode 阶段经常看到 GPU Active 很高但 SMs Active 很低——因为 kernel 太小,grid 只覆盖了少数 SM,大部分 SM 闲着。这是 batch=1 decode 的典型问题。
| 指标 | 含义 |
|---|---|
| SM Issue | SM 发射指令的比例,反映 SM 的指令吞吐 |
| Tensor Active | Tensor Core 在工作的比例。FP16 GEMM、BF16 GEMM、INT8 等都会用 Tensor Core |
Tensor Active 是 LLM 推理最关键的指标之一:
- Llama2 用 fp16/bf16 跑,GEMM 应该走 Tensor Core
- Prefill 阶段应该看到 Tensor Active 很高(GEMM 主导)
- Decode 阶段 Tensor Active 会下降(因为 GEMM 变成 GEMV 形态,效率低,而且各种 elementwise kernel 不用 Tensor Core)
- 如果 Tensor Active 一直很低,说明:① batch size 太小让 Tensor Core 吃不饱;② 走错了 kernel(用了 FP32 fallback)
这一组告诉你每个活跃 SM 里有多少 warp 在跑。一个 SM 最多承载固定数量的 warp(取决于 GPU 架构和 kernel 的资源占用)。
| 指标 | 含义 |
|---|---|
| Vertex/Tess/Geometry Warps in Flight | 图形管线相关 warp。对 CUDA compute 永远是 0,你做 LLM 推理可以无视 |
| Pixel Warps in Flight | 像素着色 warp。对 CUDA compute 也是 0 |
| Compute Warps in Flight | 真正的 CUDA compute warp。这才是你关心的 |
| Unallocated Warps in Active SMs | 活跃 SM 上还能再装多少 warp 但没装满。值越大说明 occupancy 越低 |
Warp Occupancy 是 GPU 性能优化的核心概念:
- 一个 warp 等内存时,SM 可以切到另一个 warp,这就是 latency hiding
- 如果 warp 数量不够(occupancy 低),内存延迟就藏不住,kernel 跑得慢
- 但 occupancy 不是越高越好,有些 kernel 故意用大量寄存器/shared memory 牺牲 occupancy 换算力
Llama2 推理常见情况:
- Prefill:Compute Warps 多,occupancy 高,SMs Active 高
- Decode:Compute Warps 少(grid 小),Unallocated Warps 多,SMs Active 低
| 指标 | 含义 |
|---|---|
| DRAM Read Bandwidth | 从 GPU 显存读的带宽(GB/s) |
| DRAM Write Bandwidth | 往 GPU 显存写的带宽 |
A40 的显存带宽峰值大概 696 GB/s。
关键判断:你的 kernel 是 compute-bound 还是 memory-bound?
- DRAM 带宽接近峰值(比如 > 500 GB/s on A40)→ memory-bound,优化方向是减少内存访问、提高复用
- DRAM 带宽很低 + Tensor Active 高 → compute-bound,优化方向是减少计算量
- DRAM 带宽低 + Tensor Active 低 → launch-bound 或 latency-bound,kernel 太小填不满
对 Llama2 7B 的预期:
- Decode 阶段几乎一定是 memory-bound——每个 step 要读完所有权重(7B × 2 bytes = 14GB),典型 batch=1 时 DRAM Read 接近峰值
- Prefill 阶段(长 prompt)是 compute-bound——GEMM 在算力上吃满,DRAM 带宽相对不那么紧张
这个观察(decode memory-bound vs prefill compute-bound)是所有 LLM 推理优化方案的基础。
| 指标 | 含义 |
|---|---|
| PCIe RX Throughput | GPU 通过 PCIe 接收数据的吞吐(CPU → GPU 方向) |
| PCIe TX Throughput | GPU 通过 PCIe 发送数据的吞吐(GPU → CPU 方向) |
| PCIe Read Requests to BAR1 | CPU 通过 BAR1 直接读 GPU 显存的请求数 |
| PCIe Write Requests to BAR1 | CPU 通过 BAR1 直接写 GPU 显存的请求数 |
A40 是 PCIe Gen4 x16,带宽峰值约 32 GB/s(单向)。
实战用途:
- 模型加载时会看到 PCIe RX 很高(权重从 CPU 内存搬到 GPU)
- Inference 中如果 PCIe TX 持续较高,说明频繁有数据从 GPU 回 CPU——可能是采样取 token、可能是日志输出、可能是多卡场景下的卡间通信
- 多卡 NCCL 通信走 NVLink 不走 PCIe(如果有 NVLink),但 A40 不一定有 NVLink
按优先级:
| 优先级 | 指标 | 看什么 |
|---|---|---|
| ★★★ | SMs Active | 多少 SM 在用。decode 阶段如果只有 30-50%,说明 underutilized |
| ★★★ | Tensor Active | Tensor Core 用没用。fp16/bf16 推理应该看到这个有较高的峰值 |
| ★★★ | DRAM Read Bandwidth | decode 阶段是不是 memory-bound 看这个,接近峰值就对了 |
| ★★ | GPU Active | 有没有大段 idle |
| ★★ | Unallocated Warps | occupancy 够不够 |
| ★ | GPC Clock | 有没有降频 |
| ★ | PCIe Throughput | 有没有意外的数据传输 |
如果一切正常,你应该看到:
GPC Clock: ~1.7 GHz(满频)
GPU Active: 90%+
SMs Active: 30%-60%(decode 本来就难填满)
Tensor Active: 有间歇性峰值(GEMM kernel 跑时),平均较低
Compute Warps: 不饱和,有 Unallocated Warps
DRAM Read BW: 400-600 GB/s(接近 A40 峰值 696)→ memory-bound 确诊
DRAM Write BW: 远低于 Read
PCIe TX/RX: 基本为 0(除了偶尔取 token)
这就是教科书式的 LLM decode profile——memory-bound,SM 没充分利用,但 DRAM 带宽吃满,是单 batch decode 的物理上限。
- 10 kHz 采样有抖动:每 100μs 一次的采样,对几十 μs 的小 kernel 可能漏看。看到 SMs Active 曲线尖峰跳动是正常的。
- 图形相关指标(Vertex/Pixel)对 CUDA 永远是 0:不用关心。
--gpu-metrics-device=all有性能 overhead:profile 出来的绝对时间会比实际跑稍慢,但相对关系仍然可信。- GPU Active vs SMs Active 容易混淆:前者是 GPU 整体,后者特指 SM。