1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
411
412
413
414
415
416
417
418
419
420
421
422
423
424
425
426
427
428
429
430
431
432
433
434
435
436
437
438
439
440
441
442
443
444
445
446
447
448
449
450
451
452
453
454
455
456
457
458
459
460
461
462
463
464
465
466
467
468
469
470
471
472
473
474
475
476
477
478
479
480
481
482
483
484
485
486
487
488
489
490
491
492
493
494
495
496
497
498
499
500
501
502
503
504
505
506
507
508
509
510
511
512
513
514
515
516
517
518
519
520
521
522
523
524
525
526
527
528
529
530
531
532
533
534
535
536
537
538
539
540
541
542
543
544
545
546
547
548
549
550
551
552
553
554
555
556
557
558
559
560
561
562
563
564
565
566
567
568
569
570
571
572
573
574
575
576
577
578
579
580
581
582
583
584
585
586
587
588
589
590
591
592
593
594
595
596
597
598
599
600
601
602
603
604
605
606
607
608
609
610
611
612
613
614
615
616
617
618
619
620
621
622
623
624
625
626
627
628
629
630
631
632
633
634
635
636
637
638
639
640
641
642
643
644
645
646
647
648
649
650
651
652
653
654
655
656
657
658
659
660
661
662
663
664
665
666
667
668
669
670
671
672
673
674
675
676
677
678
679
680
681
682
683
684
685
686
687
688
689
690
691
692
693
694
695
696
697
698
699
700
701
702
703
704
705
706
707
708
709
710
711
712
713
714
715
716
717
718
719
720
721
722
723
724
725
726
727
728
729
730
731
732
733
734
735
736
737
738
739
740
741
742
743
744
745
746
747
748
749
750
751
752
753
754
755
756
757
758
759
760
761
762
763
764
765
766
767
768
769
770
771
772
773
774
775
776
777
778
779
780
781
782
783
784
785
786
787
788
789
790
791
792
793
794
795
796
797
798
799
800
801
802
803
804
805
806
807
808
809
810
811
812
813
814
815
816
817
818
819
820
821
822
823
824
825
826
827
828
829
830
831
832
833
834
835
836
837
838
839
840
841
842
843
844
845
846
847
848
849
850
851
852
853
854
855
856
857
858
859
860
861
862
863
864
865
866
867
868
869
870
871
872
873
874
875
876
877
878
879
880
881
882
883
884
885
886
887
888
889
890
891
892
893
894
895
896
897
898
899
900
901
902
903
904
905
906
907
908
909
910
911
912
913
914
915
916
917
918
919
920
921
922
923
924
925
926
927
928
929
930
931
932
933
934
935
936
937
938
939
940
941
942
943
944
945
946
947
948
949
950
951
952
953
954
955
956
957
958
959
960
961
962
963
964
965
966
967
968
969
970
971
972
973
974
975
976
977
978
979
980
981
982
983
984
985
986
987
988
989
990
991
992
993
994
995
996
997
998
999
1000
1001
1002
1003
1004
1005
1006
1007
1008
1009
1010
1011
1012
1013
1014
1015
1016
1017
1018
1019
1020
1021
1022
1023
1024
1025
1026
1027
1028
1029
1030
1031
1032
1033
1034
1035
1036
1037
1038
1039
1040
1041
1042
1043
1044
1045
1046
1047
1048
1049
1050
1051
1052
1053
1054
1055
1056
1057
1058
1059
1060
1061
1062
1063
1064
1065
1066
1067
1068
1069
1070
1071
1072
1073
1074
1075
1076
1077
1078
1079
1080
1081
1082
1083
1084
1085
1086
1087
1088
1089
1090
1091
1092
1093
1094
1095
1096
1097
1098
1099
1100
1101
1102
1103
1104
1105
1106
1107
1108
1109
1110
1111
1112
1113
1114
1115
1116
1117
1118
1119
1120
1121
1122
1123
1124
1125
1126
1127
1128
1129
1130
1131
1132
1133
1134
1135
1136
1137
1138
1139
1140
1141
1142
1143
1144
1145
1146
1147
1148
1149
1150
1151
1152
1153
1154
1155
1156
1157
1158
1159
1160
1161
1162
1163
1164
1165
1166
1167
1168
1169
1170
1171
1172
1173
1174
1175
1176
1177
1178
1179
1180
1181
1182
1183
1184
1185
1186
1187
1188
1189
1190
1191
1192
1193
1194
1195
1196
1197
1198
1199
1200
1201
1202
1203
1204
1205
1206
1207
1208
1209
1210
1211
1212
1213
1214
1215
1216
1217
1218
1219
1220
1221
1222
1223
1224
1225
1226
1227
1228
1229
1230
1231
1232
1233
1234
1235
1236
1237
1238
1239
1240
1241
1242
1243
1244
1245
1246
1247
1248
1249
1250
1251
1252
1253
1254
1255
1256
1257
1258
1259
1260
1261
1262
1263
1264
1265
1266
1267
1268
1269
1270
1271
1272
1273
1274
1275
1276
1277
1278
1279
1280
1281
1282
1283
1284
1285
1286
1287
1288
1289
1290
1291
1292
1293
1294
1295
1296
1297
1298
1299
1300
1301
1302
1303
1304
1305
1306
1307
1308
1309
1310
1311
1312
1313
1314
1315
1316
1317
1318
1319
1320
1321
1322
1323
1324
1325
1326
1327
1328
1329
1330
1331
1332
1333
1334
1335
1336
1337
1338
1339
1340
1341
1342
1343
1344
1345
1346
1347
1348
1349
1350
1351
1352
1353
1354
1355
1356
1357
1358
1359
1360
1361
1362
1363
1364
1365
1366
1367
1368
1369
1370
1371
1372
1373
1374
1375
1376
1377
1378
1379
1380
1381
1382
1383
1384
1385
1386
1387
1388
1389
1390
1391
1392
1393
1394
1395
1396
1397
1398
1399
1400
1401
1402
1403
1404
1405
1406
1407
1408
1409
1410
1411
1412
1413
1414
1415
1416
1417
1418
1419
1420
1421
1422
1423
1424
1425
1426
1427
1428
1429
1430
1431
1432
1433
1434
1435
1436
1437
1438
1439
1440
1441
1442
1443
1444
1445
1446
1447
1448
1449
1450
1451
1452
1453
1454
1455
1456
1457
1458
1459
1460
1461
1462
1463
1464
1465
1466
1467
1468
1469
1470
1471
1472
1473
1474
1475
1476
1477
1478
1479
1480
1481
1482
1483
1484
1485
1486
1487
1488
1489
1490
1491
1492
1493
1494
1495
1496
1497
1498
1499
1500
1501
1502
1503
1504
1505
1506
1507
1508
1509
1510
1511
1512
1513
1514
1515
1516
1517
1518
1519
1520
1521
1522
1523
1524
1525
1526
1527
1528
1529
1530
1531
1532
1533
1534
1535
1536
1537
1538
1539
1540
1541
1542
1543
1544
1545
1546
1547
1548
1549
1550
1551
1552
1553
1554
1555
1556
1557
1558
1559
1560
1561
1562
1563
1564
1565
1566
1567
1568
1569
1570
1571
1572
1573
1574
1575
1576
1577
1578
1579
1580
1581
1582
1583
1584
1585
1586
1587
1588
1589
1590
1591
1592
1593
1594
1595
1596
1597
1598
1599
1600
1601
1602
1603
1604
1605
1606
1607
1608
1609
1610
1611
1612
1613
1614
1615
1616
1617
1618
1619
1620
1621
1622
1623
1624
1625
1626
1627
1628
1629
1630
1631
1632
1633
1634
1635
1636
1637
1638
1639
1640
1641
1642
1643
1644
1645
1646
1647
1648
1649
1650
1651
1652
1653
1654
1655
1656
1657
1658
1659
1660
1661
1662
1663
1664
1665
1666
1667
1668
1669
1670
1671
1672
1673
1674
1675
1676
1677
1678
1679
1680
1681
1682
1683
1684
1685
1686
1687
1688
1689
1690
1691
1692
1693
1694
1695
1696
1697
1698
1699
1700
1701
1702
1703
1704
1705
1706
1707
1708
1709
1710
1711
1712
1713
1714
1715
1716
1717
1718
1719
1720
1721
1722
1723
1724
1725
1726
1727
1728
1729
1730
1731
1732
1733
1734
1735
1736
1737
1738
1739
1740
1741
1742
1743
1744
1745
1746
1747
1748
1749
1750
1751
1752
1753
1754
1755
1756
1757
1758
1759
1760
1761
1762
1763
1764
1765
1766
1767
1768
1769
1770
1771
1772
1773
1774
1775
1776
1777
1778
1779
1780
1781
1782
1783
1784
1785
1786
1787
1788
1789
1790
1791
1792
1793
1794
1795
1796
1797
1798
1799
1800
1801
1802
1803
1804
1805
1806
1807
1808
1809
1810
1811
1812
1813
1814
1815
1816
1817
1818
1819
1820
1821
1822
1823
1824
1825
1826
1827
1828
1829
1830
1831
1832
1833
1834
1835
1836
1837
1838
1839
1840
1841
1842
1843
1844
1845
1846
1847
1848
1849
1850
1851
1852
1853
1854
1855
1856
1857
1858
1859
1860
1861
1862
1863
1864
1865
1866
1867
1868
1869
1870
1871
1872
1873
1874
1875
1876
1877
1878
1879
1880
1881
1882
1883
1884
1885
1886
1887
1888
1889
1890
1891
1892
1893
1894
1895
1896
1897
1898
1899
1900
1901
1902
1903
1904
1905
1906
1907
1908
1909
1910
1911
1912
1913
1914
1915
1916
1917
1918
1919
1920
1921
1922
1923
1924
1925
1926
1927
1928
1929
1930
1931
1932
1933
1934
1935
1936
1937
1938
1939
1940
1941
1942
1943
1944
1945
1946
1947
1948
1949
1950
1951
1952
1953
1954
1955
1956
1957
1958
1959
1960
1961
1962
1963
1964
1965
1966
1967
1968
1969
1970
1971
1972
1973
1974
1975
1976
1977
1978
1979
1980
1981
1982
1983
1984
1985
1986
1987
1988
1989
1990
1991
1992
1993
1994
1995
1996
1997
1998
1999
2000
2001
2002
2003
2004
2005
2006
2007
2008
2009
2010
2011
2012
2013
2014
2015
2016
2017
2018
2019
2020
2021
2022
2023
2024
2025
2026
2027
2028
2029
2030
2031
2032
2033
2034
2035
2036
2037
2038
2039
2040
2041
2042
2043
2044
2045
2046
2047
2048
2049
2050
2051
2052
2053
2054
2055
2056
2057
2058
2059
2060
2061
2062
2063
2064
2065
2066
2067
2068
2069
2070
2071
2072
2073
2074
2075
2076
2077
2078
2079
2080
2081
2082
2083
2084
2085
2086
2087
2088
2089
2090
2091
2092
2093
2094
2095
2096
2097
2098
2099
2100
2101
2102
2103
2104
2105
2106
2107
2108
2109
2110
2111
2112
2113
2114
2115
2116
2117
2118
2119
2120
2121
2122
2123
2124
2125
2126
2127
2128
2129
2130
2131
2132
2133
2134
2135
2136
2137
2138
2139
2140
2141
2142
2143
2144
2145
2146
2147
2148
2149
2150
2151
2152
2153
2154
2155
2156
2157
2158
2159
2160
2161
2162
2163
2164
2165
2166
2167
2168
2169
2170
2171
2172
2173
2174
2175
2176
2177
2178
2179
2180
2181
2182
2183
2184
2185
2186
2187
2188
2189
2190
2191
2192
2193
2194
2195
2196
2197
2198
2199
2200
2201
2202
2203
2204
2205
2206
2207
2208
2209
2210
2211
2212
2213
2214
2215
2216
2217
2218
2219
2220
2221
2222
2223
2224
2225
2226
2227
2228
2229
2230
2231
2232
2233
2234
2235
2236
2237
2238
2239
2240
2241
2242
2243
2244
2245
2246
2247
2248
2249
2250
2251
2252
2253
2254
2255
2256
2257
2258
2259
2260
2261
2262
2263
2264
2265
2266
2267
2268
2269
2270
2271
2272
2273
2274
2275
2276
2277
2278
2279
2280
2281
2282
2283
2284
2285
2286
2287
2288
2289
2290
2291
2292
2293
2294
2295
2296
2297
2298
2299
2300
2301
2302
2303
2304
2305
2306
2307
2308
2309
2310
2311
2312
2313
2314
2315
2316
2317
2318
2319
2320
2321
2322
2323
2324
2325
2326
2327
2328
2329
2330
2331
2332
2333
2334
2335
2336
2337
2338
2339
2340
2341
2342
2343
2344
2345
2346
2347
2348
2349
2350
2351
2352
2353
2354
2355
2356
2357
2358
2359
2360
2361
2362
2363
2364
2365
2366
2367
2368
2369
2370
2371
2372
2373
2374
2375
2376
2377
2378
2379
2380
2381
2382
2383
2384
2385
2386
2387
2388
2389
2390
2391
2392
2393
2394
2395
2396
2397
2398
2399
2400
2401
2402
2403
2404
2405
2406
2407
2408
2409
2410
2411
2412
2413
2414
2415
2416
2417
2418
2419
2420
2421
2422
2423
2424
2425
2426
2427
2428
2429
2430
2431
2432
2433
2434
2435
2436
2437
2438
2439
2440
2441
2442
2443
2444
2445
2446
2447
2448
2449
2450
2451
2452
2453
2454
2455
2456
2457
2458
2459
2460
2461
2462
2463
2464
2465
2466
2467
2468
2469
2470
2471
2472
2473
2474
2475
2476
2477
2478
2479
2480
2481
2482
2483
2484
2485
2486
2487
2488
2489
2490
2491
2492
2493
2494
2495
2496
2497
2498
2499
2500
2501
2502
2503
2504
2505
2506
2507
2508
2509
2510
2511
2512
2513
2514
2515
2516
2517
2518
2519
2520
2521
2522
2523
2524
2525
2526
2527
2528
2529
2530
2531
2532
2533
2534
2535
2536
2537
2538
2539
2540
2541
2542
|
<!DOCTYPE html>
<html class="writer-html5" lang="en" >
<head>
<meta charset="utf-8" /><meta name="generator" content="Docutils 0.17.1: http://docutils.sourceforge.net/" />
<meta content="The User Guide for Nsight Compute Profiling Guide." name="description" />
<meta content="User Guide" name="keywords" />
<meta name="viewport" content="width=device-width, initial-scale=1.0" />
<title>2. Kernel Profiling Guide — NsightCompute 12.4 documentation</title>
<link rel="stylesheet" href="../_static/pygments.css" type="text/css" />
<link rel="stylesheet" href="../_static/css/theme.css" type="text/css" />
<link rel="stylesheet" href="../_static/design-style.b7bb847fb20b106c3d81b95245e65545.min.css" type="text/css" />
<link rel="stylesheet" href="../_static/omni-style.css" type="text/css" />
<link rel="stylesheet" href="../_static/api-styles.css" type="text/css" />
<link rel="shortcut icon" href="../_static/nsight-compute.ico"/>
<!--[if lt IE 9]>
<script src="../_static/js/html5shiv.min.js"></script>
<![endif]-->
<script data-url_root="../" id="documentation_options" src="../_static/documentation_options.js"></script>
<script src="../_static/jquery.js"></script>
<script src="../_static/underscore.js"></script>
<script src="../_static/doctools.js"></script>
<script src="../_static/mermaid-init.js"></script>
<script src="../_static/design-tabs.js"></script>
<script src="../_static/version.js"></script>
<script src="../_static/social-media.js"></script>
<script src="../_static/js/theme.js"></script>
<link rel="index" title="Index" href="../genindex.html" />
<link rel="search" title="Search" href="../search.html" />
<link rel="next" title="3. Nsight Compute" href="../NsightCompute/index.html" />
<link rel="prev" title="1. Release Notes" href="../ReleaseNotes/index.html" />
</head>
<body class="wy-body-for-nav">
<div class="wy-grid-for-nav">
<nav data-toggle="wy-nav-shift" class="wy-nav-side">
<div class="wy-side-scroll">
<div class="wy-side-nav-search" >
<a href="../index.html">
<img src="../_static/nsight-compute.png" class="logo" alt="Logo"/>
</a>
<div role="search">
<form id="rtd-search-form" class="wy-form" action="../search.html" method="get">
<input type="text" name="q" placeholder="Search docs" />
<input type="hidden" name="check_keywords" value="yes" />
<input type="hidden" name="area" value="default" />
</form>
</div>
</div><div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="Navigation menu">
<p class="caption" role="heading"><span class="caption-text">Nsight Compute</span></p>
<ul class="current">
<li class="toctree-l1"><a class="reference internal" href="../ReleaseNotes/index.html">1. Release Notes</a></li>
<li class="toctree-l1 current"><a class="current reference internal" href="#">2. Kernel Profiling Guide</a><ul>
<li class="toctree-l2"><a class="reference internal" href="#introduction">2.1. Introduction</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#profiling-applications">2.1.1. Profiling Applications</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#metric-collection">2.2. Metric Collection</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#sets-and-sections">2.2.1. Sets and Sections</a></li>
<li class="toctree-l3"><a class="reference internal" href="#sections-and-rules">2.2.2. Sections and Rules</a></li>
<li class="toctree-l3"><a class="reference internal" href="#replay">2.2.3. Replay</a><ul>
<li class="toctree-l4"><a class="reference internal" href="#kernel-replay">Kernel Replay</a></li>
<li class="toctree-l4"><a class="reference internal" href="#application-replay">Application Replay</a></li>
<li class="toctree-l4"><a class="reference internal" href="#range-replay">Range Replay</a><ul>
<li class="toctree-l5"><a class="reference internal" href="#defining-ranges">Defining Ranges</a></li>
<li class="toctree-l5"><a class="reference internal" href="#supported-apis">Supported APIs</a></li>
</ul>
</li>
<li class="toctree-l4"><a class="reference internal" href="#application-range-replay">Application Range Replay</a></li>
<li class="toctree-l4"><a class="reference internal" href="#graph-profiling">Graph Profiling</a></li>
</ul>
</li>
<li class="toctree-l3"><a class="reference internal" href="#compatibility">2.2.4. Compatibility</a></li>
<li class="toctree-l3"><a class="reference internal" href="#profile-series">2.2.5. Profile Series</a></li>
<li class="toctree-l3"><a class="reference internal" href="#overhead">2.2.6. Overhead</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#metrics-guide">2.3. Metrics Guide</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#hardware-model">2.3.1. Hardware Model</a></li>
<li class="toctree-l3"><a class="reference internal" href="#metrics-structure">2.3.2. Metrics Structure</a></li>
<li class="toctree-l3"><a class="reference internal" href="#metrics-decoder">2.3.3. Metrics Decoder</a></li>
<li class="toctree-l3"><a class="reference internal" href="#range-and-precision">2.3.4. Range and Precision</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#metrics-reference">2.4. Metrics Reference</a></li>
<li class="toctree-l2"><a class="reference internal" href="#sampling">2.5. Sampling</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#pm-sampling">2.5.1. PM Sampling</a></li>
<li class="toctree-l3"><a class="reference internal" href="#warp-sampling">2.5.2. Warp Sampling</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#reproducibility">2.6. Reproducibility</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#serialization">2.6.1. Serialization</a></li>
<li class="toctree-l3"><a class="reference internal" href="#clock-control">2.6.2. Clock Control</a></li>
<li class="toctree-l3"><a class="reference internal" href="#cache-control">2.6.3. Cache Control</a></li>
<li class="toctree-l3"><a class="reference internal" href="#persistence-mode">2.6.4. Persistence Mode</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#special-configurations">2.7. Special Configurations</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#multi-instance-gpu">2.7.1. Multi Instance GPU</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#roofline-charts">2.8. Roofline Charts</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#roofline-overview">2.8.1. Overview</a></li>
<li class="toctree-l3"><a class="reference internal" href="#analysis">2.8.2. Analysis</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#memory-chart">2.9. Memory Chart</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#memory-chart-overview">2.9.1. Overview</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#memory-tables">2.10. Memory Tables</a><ul>
<li class="toctree-l3"><a class="reference internal" href="#shared-memory">2.10.1. Shared Memory</a></li>
<li class="toctree-l3"><a class="reference internal" href="#l1-tex-cache">2.10.2. L1/TEX Cache</a></li>
<li class="toctree-l3"><a class="reference internal" href="#l2-cache">2.10.3. L2 Cache</a></li>
<li class="toctree-l3"><a class="reference internal" href="#l2-cache-eviction-policies">2.10.4. L2 Cache Eviction Policies</a></li>
<li class="toctree-l3"><a class="reference internal" href="#device-memory">2.10.5. Device Memory</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="#faq">2.11. FAQ</a></li>
</ul>
</li>
<li class="toctree-l1"><a class="reference internal" href="../NsightCompute/index.html">3. Nsight Compute</a></li>
<li class="toctree-l1"><a class="reference internal" href="../NsightComputeCli/index.html">4. Nsight Compute CLI</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">Developer Interfaces</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../CustomizationGuide/index.html">1. Customization Guide</a></li>
<li class="toctree-l1"><a class="reference internal" href="../NvRulesAPI/index.html">2. NvRules API</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">Training</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../Training/index.html">Training</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">Release Information</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../Archives/index.html">Archives</a></li>
</ul>
<p class="caption" role="heading"><span class="caption-text">Copyright and Licenses</span></p>
<ul>
<li class="toctree-l1"><a class="reference internal" href="../CopyrightAndLicenses/index.html">Copyright and Licenses</a></li>
</ul>
</div>
</div>
</nav>
<section data-toggle="wy-nav-shift" class="wy-nav-content-wrap"><nav class="wy-nav-top" aria-label="Mobile navigation menu" >
<i data-toggle="wy-nav-top" class="fa fa-bars"></i>
<a href="../index.html">NsightCompute</a>
</nav>
<div class="wy-nav-content">
<div class="rst-content">
<div role="navigation" aria-label="Page navigation">
<ul class="wy-breadcrumbs">
<li><a href="../index.html" class="icon icon-home"></a> »</li>
<li><span class="section-number">2. </span>Kernel Profiling Guide</li>
<li class="wy-breadcrumbs-aside">
</li>
<li class="wy-breadcrumbs-aside">
<span>v2024.1.1 |</span>
<a href="https://developer.nvidia.com/nsight-compute-history" class="reference external">Archive</a>
<span> </span>
</li>
</ul>
<hr/>
</div>
<div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
<div itemprop="articleBody">
<section id="kernel-profiling-guide">
<h1><span class="section-number">2. </span>Kernel Profiling Guide<a class="headerlink" href="#kernel-profiling-guide" title="Permalink to this headline"></a></h1>
<p>Nsight Compute profiling guide.</p>
<p>Kernel Profiling Guide with metric types and meaning, data collection modes and FAQ for common problems.</p>
<section id="introduction">
<h2><span class="section-number">2.1. </span>Introduction<a class="headerlink" href="#introduction" title="Permalink to this headline"></a></h2>
<p>This guide describes various profiling topics related to NVIDIA Nsight Compute and NVIDIA Nsight Compute CLI. Most of these apply to both the UI and the CLI version of the tool.</p>
<p>To use the tools effectively, it is recommended to read this guide, as well as at least the following chapters of the <em>CUDA Programming Guide</em>:</p>
<ul class="simple">
<li><p><a class="reference external" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programming-model">Programming Model</a></p></li>
<li><p><a class="reference external" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#hardware-implementation">Hardware Implementation</a></p></li>
<li><p><a class="reference external" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#performance-guidelines">Performance Guidelines</a></p></li>
</ul>
<p>Afterwards, it should be enough to read the <em>Quickstart</em> chapter of the NVIDIA Nsight Compute or NVIDIA Nsight Compute CLI documentation, respectively, to start using the tools.</p>
<section id="profiling-applications">
<h3><span class="section-number">2.1.1. </span>Profiling Applications<a class="headerlink" href="#profiling-applications" title="Permalink to this headline"></a></h3>
<p>During regular execution, a CUDA application process will be launched by the user. It communicates directly with the CUDA user-mode driver, and potentially with the CUDA runtime library.</p>
<figure class="align-center" id="id8">
<img alt="../_images/regular-application-process.png" src="../_images/regular-application-process.png" />
<figcaption>
<p><span class="caption-text">Regular Application Execution</span><a class="headerlink" href="#id8" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p>When profiling an application with NVIDIA Nsight Compute, the behavior is different. The user launches the NVIDIA Nsight Compute frontend (either the UI or the CLI) on the host system, which in turn starts the actual application as a new process on the target system. While host and target are often the same machine, the target can also be a remote system with a potentially different operating system.</p>
<p>The tool inserts its measurement libraries into the application process, which allow the profiler to intercept communication with the CUDA user-mode driver. In addition, when a kernel launch is detected, the libraries can collect the requested performance metrics from the GPU. The results are then transferred back to the frontend.</p>
<figure class="align-center" id="id9">
<img alt="../_images/profiled-process.png" src="../_images/profiled-process.png" />
<figcaption>
<p><span class="caption-text">Profiled Application Execution</span><a class="headerlink" href="#id9" title="Permalink to this image"></a></p>
</figcaption>
</figure>
</section>
</section>
<section id="metric-collection">
<h2><span class="section-number">2.2. </span>Metric Collection<a class="headerlink" href="#metric-collection" title="Permalink to this headline"></a></h2>
<p>Collection of performance metrics is the key feature of NVIDIA Nsight Compute. Since there is a huge list of metrics available, it is often easier to use some of the tool’s pre-defined <a class="reference external" href="index.html#sets-and-sections">sets or sections</a> to collect a commonly used subset. Users are free to adjust which metrics are collected for which kernels as needed, but it is important to keep in mind the <a class="reference external" href="index.html#overhead">Overhead</a> associated with data collection.</p>
<section id="sets-and-sections">
<h3><span class="section-number">2.2.1. </span>Sets and Sections<a class="headerlink" href="#sets-and-sections" title="Permalink to this headline"></a></h3>
<p>NVIDIA Nsight Compute uses <em>Section Sets</em> (short <em>sets</em>) to decide, on a very high level, the amount of metrics to be collected. Each set includes one or more <em>Sections</em>, with each section specifying several logically associated metrics. For example, one section might include only high-level SM and memory utilization metrics, while another could include metrics associated with the memory units, or the HW scheduler.</p>
<p>The number and type of metrics specified by a section has significant impact on the overhead during profiling. To allow you to quickly choose between a fast, less detailed profile and a slower, more comprehensive analysis, you can select the respective section set. See <a class="reference external" href="index.html#overhead">Overhead</a> for more information on profiling overhead.</p>
<p>By default, a relatively small number of metrics is collected. Those mostly include high-level utilization information as well as static launch and occupancy data. The latter two are regularly available without replaying the kernel launch. The <code class="docutils literal notranslate"><span class="pre">basic</span></code> set is collected when no <code class="docutils literal notranslate"><span class="pre">--set</span></code>, <code class="docutils literal notranslate"><span class="pre">--section</span></code> and no <code class="docutils literal notranslate"><span class="pre">--metrics</span></code> options are passed on the command line. The full set of sections can be collected with <code class="docutils literal notranslate"><span class="pre">--set</span> <span class="pre">full</span></code>.</p>
<p>Use <code class="docutils literal notranslate"><span class="pre">--list-sets</span></code> to see the list of currently available sets. Use <code class="docutils literal notranslate"><span class="pre">--list-sections</span></code> to see the list of currently available sections. The default search directory and the location of pre-defined section files are also called <code class="docutils literal notranslate"><span class="pre">sections/</span></code>. All related command line options can be found in the NVIDIA Nsight Compute CLI documentation.</p>
</section>
<section id="sections-and-rules">
<h3><span class="section-number">2.2.2. </span>Sections and Rules<a class="headerlink" href="#sections-and-rules" title="Permalink to this headline"></a></h3>
<table class="table-no-stripes docutils align-default" id="id10">
<caption><span class="caption-text">Table 1. Available Sections</span><a class="headerlink" href="#id10" title="Permalink to this table"></a></caption>
<colgroup>
<col style="width: 12%" />
<col style="width: 88%" />
</colgroup>
<thead>
<tr class="row-odd"><th class="head"><p>Identifier and Filename</p></th>
<th class="head"><p>Description</p></th>
</tr>
</thead>
<tbody>
<tr class="row-even"><td><p>ComputeWorkloadAnalysis (Compute Workload Analysis)</p></td>
<td><p>Detailed analysis of the compute resources of the streaming multiprocessors (SM), including the achieved instructions per clock (IPC) and the utilization of each available pipeline. Pipelines with very high utilization might limit the overall performance.</p></td>
</tr>
<tr class="row-odd"><td><p>InstructionStats (Instruction Statistics)</p></td>
<td><p>Statistics of the executed low-level assembly instructions (SASS). The instruction mix provides insight into the types and frequency of the executed instructions. A narrow mix of instruction types implies a dependency on few instruction pipelines, while others remain unused. Using multiple pipelines allows hiding latencies and enables parallel execution.</p></td>
</tr>
<tr class="row-even"><td><p>LaunchStats (Launch Statistics)</p></td>
<td><p>Summary of the configuration used to launch the kernel. The launch configuration defines the size of the kernel grid, the division of the grid into blocks, and the GPU resources needed to execute the kernel. Choosing an efficient launch configuration maximizes device utilization.</p></td>
</tr>
<tr class="row-odd"><td><p>MemoryWorkloadAnalysis (Memory Workload Analysis)</p></td>
<td><p>Detailed analysis of the memory resources of the GPU. Memory can become a limiting factor for the overall kernel performance when fully utilizing the involved hardware units (Mem Busy), exhausting the available communication bandwidth between those units (Max Bandwidth), or by reaching the maximum throughput of issuing memory instructions (Mem Pipes Busy).
Depending on the limiting factor, the memory chart and tables allow to identify the exact bottleneck in the memory system.</p></td>
</tr>
<tr class="row-even"><td><p>NUMA Affinity (NumaAffinity)</p></td>
<td><p>Non-uniform memory access (NUMA) affinities based on compute and memory distances for all GPUs.</p></td>
</tr>
<tr class="row-odd"><td><p>Nvlink (Nvlink)</p></td>
<td><p>High-level summary of NVLink utilization. It shows the total received and transmitted (sent) memory, as well as the overall link peak utilization.</p></td>
</tr>
<tr class="row-even"><td><p>Nvlink_Tables (Nvlink_Tables)</p></td>
<td><p>Detailed tables with properties for each NVLink.</p></td>
</tr>
<tr class="row-odd"><td><p>Nvlink_Topology (Nvlink_Topology)</p></td>
<td><p>NVLink Topology diagram shows logical NVLink connections with transmit/receive throughput.</p></td>
</tr>
<tr class="row-even"><td><p>Occupancy (Occupancy)</p></td>
<td><p>Occupancy is the ratio of the number of active warps per multiprocessor to the maximum number of possible active warps. Another way to view occupancy is the percentage of the hardware’s ability to process warps that is actively in use.
Higher occupancy does not always result in higher performance, however, low occupancy always reduces the ability to hide latencies, resulting in overall performance degradation. Large discrepancies between the theoretical and the achieved occupancy during execution typically indicates highly imbalanced workloads.</p></td>
</tr>
<tr class="row-odd"><td><p>PM Sampling (PmSampling)</p></td>
<td><p>Timeline view of metrics sampled periodically over the workload duration. Data is collected across multiple passes. Use this section to understand how workload behavior changes over its runtime.</p></td>
</tr>
<tr class="row-even"><td><p>PM Sampling: Warp States (PmSampling_WarpStates)</p></td>
<td><p>Warp states sampled periodically over the workload duration. Metrics in different groups come from different passes.</p></td>
</tr>
<tr class="row-odd"><td><p>SchedulerStats (Scheduler Statistics)</p></td>
<td><p>Summary of the activity of the schedulers issuing instructions. Each scheduler maintains a pool of warps that it can issue instructions for. The upper bound of warps in the pool (Theoretical Warps) is limited by the launch configuration. On every cycle each scheduler checks the state of the allocated warps in the pool (Active Warps).
Active warps that are not stalled (Eligible Warps) are ready to issue their next instruction. From the set of eligible warps, the scheduler selects a single warp from which to issue one or more instructions (Issued Warp). On cycles with no eligible warps, the issue slot is skipped and no instruction is issued. Having many skipped issue slots indicates poor latency hiding.</p></td>
</tr>
<tr class="row-even"><td><p>SourceCounters (Source Counters)</p></td>
<td><p>Source metrics, including branch efficiency and sampled warp stall reasons. Warp Stall Sampling metrics are periodically sampled over the kernel runtime. They indicate when warps were stalled and couldn’t be scheduled. See the documentation for a description of all stall reasons. Only focus on stalls if the schedulers fail to issue every cycle.</p></td>
</tr>
<tr class="row-odd"><td><p>SpeedOfLight (GPU Speed Of Light Throughput)</p></td>
<td><p>High-level overview of the throughput for compute and memory resources of the GPU. For each unit, the throughput reports the achieved percentage of utilization with respect to the theoretical maximum. Breakdowns show the throughput for each individual sub-metric of Compute and Memory to clearly identify the highest contributor.</p></td>
</tr>
<tr class="row-even"><td><p>WarpStateStats (Warp State Statistics)</p></td>
<td><p>Analysis of the states in which all warps spent cycles during the kernel execution. The warp states describe a warp’s readiness or inability to issue its next instruction. The warp cycles per instruction define the latency between two consecutive instructions. The higher the value, the more warp parallelism is required to hide this latency.
For each warp state, the chart shows the average number of cycles spent in that state per issued instruction. Stalls are not always impacting the overall performance nor are they completely avoidable. Only focus on stall reasons if the schedulers fail to issue every cycle.</p></td>
</tr>
</tbody>
</table>
</section>
<section id="replay">
<h3><span class="section-number">2.2.3. </span>Replay<a class="headerlink" href="#replay" title="Permalink to this headline"></a></h3>
<p>Depending on which metrics are to be collected, kernels might need to be <em>replayed</em> one or more times, since not all metrics can be collected in a single <em>pass</em>. For example, the number of metrics originating from hardware (HW) performance counters that the GPU can collect at the same time is limited. In addition, patch-based software (SW) performance counters can have a high impact on kernel runtime and would skew results for HW counters.</p>
<section id="kernel-replay">
<h4>Kernel Replay<a class="headerlink" href="#kernel-replay" title="Permalink to this headline"></a></h4>
<p>In <em>Kernel Replay</em>, all metrics requested for a specific kernel instance in NVIDIA Nsight Compute are grouped into one or more passes. For the first pass, all GPU memory that can be accessed by the kernel is saved. After the first pass, the subset of memory that is written by the kernel is determined. Before each pass (except the first one), this subset is restored in its original location to have the kernel access the same memory contents in each replay pass.</p>
<p>NVIDIA Nsight Compute attempts to use the fastest available storage location for this save-and-restore strategy. For example, if data is allocated in device memory, and there is still enough device memory available, it is stored there directly. If it runs out of device memory, the data is transferred to the CPU host memory. Likewise, if an allocation originates from CPU host memory, the tool first attempts to save it into the same memory location, if possible.</p>
<p>As explained in <a class="reference external" href="index.html#overhead">Overhead</a>, the time needed for this increases the more memory is accessed, especially written, by a kernel. If NVIDIA Nsight Compute determines that only a single replay pass is necessary to collect the requested metrics, no save-and-restore is performed at all to reduce overhead.</p>
<figure class="align-center" id="id11">
<img alt="../_images/replay-regular-execution.png" src="../_images/replay-regular-execution.png" />
<figcaption>
<p><span class="caption-text">Regular Application Execution</span><a class="headerlink" href="#id11" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<figure class="align-center" id="id12">
<img alt="../_images/replay-kernel.png" src="../_images/replay-kernel.png" />
<figcaption>
<p><span class="caption-text">Execution with Kernel Replay. All memory is saved, and memory written by the kernel is restored in-between replay passes.</span><a class="headerlink" href="#id12" title="Permalink to this image"></a></p>
</figcaption>
</figure>
</section>
<section id="application-replay">
<h4>Application Replay<a class="headerlink" href="#application-replay" title="Permalink to this headline"></a></h4>
<p>In <em>Application Replay</em>, all metrics requested for a specific kernel launch in NVIDIA Nsight Compute are grouped into one or more passes. In contrast to <a class="reference external" href="index.html#kernel-replay">Kernel Replay</a>, the complete application is run multiple times, so that in each run one of those passes can be collected per kernel.</p>
<p>For correctly identifying and combining performance counters collected from multiple application replay passes of a single kernel launch into one result, the application needs to be deterministic with respect to its kernel activities and their assignment to GPUs, contexts, streams, and potentially NVTX ranges. Normally, this also implies that the application needs to be deterministic with respect to its overall execution.</p>
<p>Application replay has the benefit that memory accessed by the kernel does not need to be saved and restored via the tool, as each kernel launch executes only once during the lifetime of the application process. Besides avoiding memory save-and-restore overhead, application replay also allows to disable <a class="reference external" href="index.html#cache-control">Cache Control</a>. This is especially useful if other GPU activities preceding a specific kernel launch are used by the application to set caches to some expected state.</p>
<p>In addition, application replay can support profiling kernels that have interdependencies to the host during execution. With kernel replay, this class of kernels typically hangs when being profiled, because the necessary responses from the host are missing in all but the first pass. In contrast, application replay ensures the correct behavior of the program execution in each pass.</p>
<p>In contrast to kernel replay, multiple passes collected via application replay imply that all host-side activities of the application are duplicated, too. If the application requires significant time for e.g. setup or file-system access, the overhead will increase accordingly.</p>
<figure class="align-center" id="id13">
<img alt="../_images/replay-regular-execution.png" src="../_images/replay-regular-execution.png" />
<figcaption>
<p><span class="caption-text">Regular Application Execution</span><a class="headerlink" href="#id13" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<figure class="align-center" id="id14">
<img alt="../_images/replay-application.png" src="../_images/replay-application.png" />
<figcaption>
<p><span class="caption-text">Execution with Application Replay. No memory is saved or restored, but the cost of running the application itself is duplicated.</span><a class="headerlink" href="#id14" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p>Across application replay passes, NVIDIA Nsight Compute matches metric data for the individual, selected kernel launches. The matching strategy can be selected using the <code class="docutils literal notranslate"><span class="pre">--app-replay-match</span></code> option. For matching, only kernels within the same process and running on the same device are considered. By default, the <em>grid</em> strategy is used, which matches launches according to their kernel name and grid size. When multiple launches have the same attributes (e.g. name and grid size), they are matched in execution order.</p>
<figure class="align-center" id="id15">
<img alt="../_images/replay-application-kernel-matching.png" src="../_images/replay-application-kernel-matching.png" />
<figcaption>
<p><span class="caption-text">Kernel matching during application replay using the <em>grid</em> strategy.</span><a class="headerlink" href="#id15" title="Permalink to this image"></a></p>
</figcaption>
</figure>
</section>
<section id="range-replay">
<h4>Range Replay<a class="headerlink" href="#range-replay" title="Permalink to this headline"></a></h4>
<p>In <em>Range Replay</em>, all requested metrics in NVIDIA Nsight Compute are grouped into one or more passes. In contrast to <a class="reference external" href="index.html#kernel-replay">Kernel Replay</a> and <a class="reference external" href="index.html#application-replay">Application Replay</a>, <em>Range Replay</em> captures and replays complete ranges of CUDA API calls and kernel launches within the profiled application. Metrics are then not associated with individual kernels but with the entire range. This allows the tool to execute kernels without serialization and thereby supports profiling kernels that should be run concurrently for correctness or performance reasons.</p>
<figure class="align-center" id="id16">
<img alt="../_images/replay-range.png" src="../_images/replay-range.png" />
<figcaption>
<p><span class="caption-text">Execution with Range Replay. An entire range of API calls and kernel launches is captured and replayed. Host and device memory is saved and restored as necessary.</span><a class="headerlink" href="#id16" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<section id="defining-ranges">
<span id="range-replay-define-range"></span><h5>Defining Ranges<a class="headerlink" href="#defining-ranges" title="Permalink to this headline"></a></h5>
<p>Range replay requires you to specify the range for profiling in the application. A range is defined by a start and an end marker and includes all CUDA API calls and kernels launched between these markers from any CPU thread. The application is responsible for inserting appropriate synchronization between threads to ensure that the anticipated set of API calls is captured. Range markers can be set using one of the following options:</p>
<ul>
<li><p><strong>Profiler Start/Stop API</strong></p>
<p>Set the start marker using <code class="docutils literal notranslate"><span class="pre">cu(da)ProfilerStart</span></code> and the end marker using <code class="docutils literal notranslate"><span class="pre">cu(da)ProfilerStop</span></code>. Note: The CUDA driver API variants of this API require to include <code class="docutils literal notranslate"><span class="pre">cudaProfiler.h</span></code>. The CUDA runtime variants require to include <code class="docutils literal notranslate"><span class="pre">cuda_profiler_api.h</span></code>.</p>
<p>This is the default for NVIDIA Nsight Compute.</p>
</li>
<li><p><strong>NVTX Ranges</strong></p>
<p>Define the range using an <a class="reference external" href="../NsightComputeCli/index.html#nvtx-filtering">NVTX Include</a> expression. The range capture starts with the first CUDA API call and ends at the last API call for which the expression is matched, respectively. If multiple expressions are specified, a range is defined as soon as any of them matches. Hence, multiple expressions can be used to conveniently capture and profile multiple ranges for the same application execution.</p>
<p>The application must have been instrumented with the NVTX API for any expressions to match.</p>
<p>This mode is enabled by passing <code class="docutils literal notranslate"><span class="pre">--nvtx</span> <span class="pre">--nvtx-include</span> <span class="pre"><expression></span> <span class="pre">[--nvtx-include</span> <span class="pre"><expression>]</span></code> to the NVIDIA Nsight Compute CLI.</p>
</li>
</ul>
<p>Ranges must fulfill several requirements:</p>
<ul class="simple">
<li><p>It must be possible to synchronize all active CUDA contexts at the start of the range.</p></li>
<li><p>Ranges must not include unsupported CUDA API calls. See <a class="reference external" href="index.html#range-replay-supported-apis">Supported APIs</a> for the list of currently supported APIs.</p></li>
</ul>
<p>In addition, there are several recommendations that ranges should comply with to guarantee a correct capture and replay:</p>
<ul class="simple">
<li><p>Set ranges as narrow as possible for capturing a specific set of CUDA kernel lanuches. The more API calls are included, the higher the potentially created overhead from capturing and replaying these API calls.</p></li>
<li><p>Avoid freeing host allocations written by device memory during the range. This includes both heap as well as stack allocations. NVIDIA Nsight Compute does not intercept creation or destruction of generic host (CPU)-based allocations. However, to guarantee correct program execution after any replay of the range, the tool attempts to restore host allocations that were written from device memory during the capture. If these host addresses are invalid or re-assigned, the program behavior is undefined and potentially unstable. In cases where avoiding freeing such allocations is not possible, you should limit profiling to one range using <code class="docutils literal notranslate"><span class="pre">--launch-count</span> <span class="pre">1</span></code>, set the <em>disable-host-restore</em> range replay option and optionally use <code class="docutils literal notranslate"><span class="pre">--kill</span> <span class="pre">yes</span></code> to terminate the process after this range.</p></li>
<li><p>When defining the range markers using <code class="docutils literal notranslate"><span class="pre">cu(da)ProfilerStart/Stop</span></code>, prefer the CUDA driver API calls <code class="docutils literal notranslate"><span class="pre">cuProfilerStart/Stop</span></code>. Internally, NVIDIA Nsight Compute only intercepts the CUDA driver API variants and the CUDA runtime API may not trigger these if no CUDA context is active on the calling thread.</p></li>
</ul>
</section>
<section id="supported-apis">
<span id="range-replay-supported-apis"></span><h5>Supported APIs<a class="headerlink" href="#supported-apis" title="Permalink to this headline"></a></h5>
<p>Range replay supports a subset of the CUDA API for capture and replay. This page lists the supported functions as well as any further, API-specific limitations that may apply. If an unsupported API call is detected in the captured range, an error is reported and the range cannot be profiled. The groups listed below match the ones found in the <a class="reference external" href="https://docs.nvidia.com/cuda/cuda-driver-api/index.html">CUDA Driver API documentation</a>.</p>
<p>Generally, range replay only captures and replay CUDA <em>Driver</em> API calls. CUDA <em>Runtime</em> APIs calls can be captured when they generate only supported CUDA Driver API calls internally. Deprecated APIs are not supported.</p>
<p><strong>Error Handling</strong></p>
<p>All supported.</p>
<p><strong>Initialization</strong></p>
<p>Not supported.</p>
<p><strong>Version Management</strong></p>
<p>All supported.</p>
<p><strong>Device Management</strong></p>
<p>All supported, except:</p>
<ul class="simple">
<li><p>cuDeviceSetMemPool</p></li>
</ul>
<p><strong>Primary Context Management</strong></p>
<ul class="simple">
<li><p>cuDevicePrimaryCtxGetState</p></li>
</ul>
<p><strong>Context Management</strong></p>
<p>All supported, except:</p>
<ul class="simple">
<li><p>cuCtxSetCacheConfig</p></li>
<li><p>cuCtxSetSharedMemConfig</p></li>
</ul>
<p><strong>Module Management</strong></p>
<ul class="simple">
<li><p>cuModuleGetFunction</p></li>
<li><p>cuModuleGetGlobal</p></li>
<li><p>cuModuleGetSurfRef</p></li>
<li><p>cuModuleGetTexRef</p></li>
<li><p>cuModuleLoad</p></li>
<li><p>cuModuleLoadData</p></li>
<li><p>cuModuleLoadDataEx</p></li>
<li><p>cuModuleLoadFatBinary</p></li>
<li><p>cuModuleUnload</p></li>
</ul>
<p><strong>Library Management</strong></p>
<p>All supported, except:</p>
<ul class="simple">
<li><p>cuKernelSetAttribute</p></li>
<li><p>cuKernelSetCacheConfig</p></li>
</ul>
<p><strong>Memory Management</strong></p>
<ul class="simple">
<li><p>cuArray*</p></li>
<li><p>cuDeviceGetByPCIBusId</p></li>
<li><p>cuDeviceGetPCIBusId</p></li>
<li><p>cuMemAlloc</p></li>
<li><p>cuMemAllocHost</p></li>
<li><p>cuMemAllocPitch</p></li>
<li><p>cuMemcpy*</p></li>
<li><p>cuMemFree</p></li>
<li><p>cuMemFreeHost</p></li>
<li><p>cuMemGetAddressRange</p></li>
<li><p>cuMemGetInfo</p></li>
<li><p>cuMemHostAlloc</p></li>
<li><p>cuMemHostGetDevicePointer</p></li>
<li><p>cuMemHostGetFlags</p></li>
<li><p>cuMemHostRegister</p></li>
<li><p>cuMemHostUnregister</p></li>
<li><p>cuMemset*</p></li>
<li><p>cuMipmapped*</p></li>
</ul>
<p><strong>Virtual Memory Management</strong></p>
<p>Not supported.</p>
<p><strong>Stream Ordered Memory Allocator</strong></p>
<p>Not supported.</p>
<p><strong>Unified Addressing</strong></p>
<p>Not supported.</p>
<p><strong>Stream Management</strong></p>
<ul class="simple">
<li><p>cuStreamCreate*</p></li>
<li><p>cuStreamDestroy</p></li>
<li><p>cuStreamGet*</p></li>
<li><p>cuStreamQuery</p></li>
<li><p>cuStreamSetAttribute</p></li>
<li><p>cuStreamSynchronize</p></li>
<li><p>cuStreamWaitEvent</p></li>
</ul>
<p><strong>Event Management</strong></p>
<p>All supported.</p>
<p><strong>External Resource interoperability</strong></p>
<p>Not supported.</p>
<p><strong>Stream Memory Operations</strong></p>
<p>Not supported.</p>
<p><strong>Execution Control</strong></p>
<ul class="simple">
<li><p>cuFuncGetAttribute</p></li>
<li><p>cuFuncGetModule</p></li>
<li><p>cuFuncSetAttribute</p></li>
<li><p>cuFuncSetCacheConfig</p></li>
<li><p>cuLaunchCooperativeKernel</p></li>
<li><p>cuLaunchHostFunc</p></li>
<li><p>cuLaunchKernel</p></li>
</ul>
<p><strong>Graph Management</strong></p>
<p>Not supported.</p>
<p><strong>Occupancy</strong></p>
<p>All supported.</p>
<p><strong>Texture/Surface Reference Management</strong></p>
<p>Not supported.</p>
<p><strong>Texture Object Management</strong></p>
<p>All supported.</p>
<p><strong>Surface Object Management</strong></p>
<p>All supported.</p>
<p><strong>Peer Context Memory Access</strong></p>
<p>Not supported.</p>
<p><strong>Graphics Interoperability</strong></p>
<p>Not supported.</p>
<p><strong>Driver Entry Point Access</strong></p>
<p>All supported.</p>
<p><strong>Surface Object Management</strong></p>
<p>All supported.</p>
<p><strong>OpenGL Interoperability</strong></p>
<p>Not supported.</p>
<p><strong>VDPAU Interoperability</strong></p>
<p>Not supported.</p>
<p><strong>EGL Interoperability</strong></p>
<p>Not supported.</p>
<p><strong>Green Contexts</strong></p>
<ul class="simple">
<li><p>cuCtxFromGreenCtx</p></li>
<li><p>cuGreenCtxCreate</p></li>
<li><p>cuGreenCtxDestroy</p></li>
<li><p>cuGreenCtxRecordEvent</p></li>
<li><p>cuGreenCtxWaitEvent</p></li>
<li><p>cuStreamGetGreenCtx</p></li>
</ul>
</section>
</section>
<section id="application-range-replay">
<h4>Application Range Replay<a class="headerlink" href="#application-range-replay" title="Permalink to this headline"></a></h4>
<p>In <em>Application Range Replay</em>, all requested metrics in NVIDIA Nsight Compute are grouped into one or more passes. Similar to <a class="reference external" href="index.html#range-replay">Range Replay</a>, metrics are not associated with individual kernels but with the entire selected range. This allows the tool to execute workloads (kernels, CUDA graphs, …) without serialization and thereby supports profiling workloads that must be run concurrently for correctness or performance reasons.</p>
<p>In contrast to Range Replay, the range is not explicitly captured and executed directly for each pass, but instead the entire application is re-run multiple times, with one pass collected for each range in every application execution. This has the benefit that no application state must be observed and captured for each range and API calls within the range do not need to be supported explicitly, as correct execution of the range is handled by the application itself.</p>
<p>Defining ranges to profile is identical to <a class="reference external" href="index.html#range-replay-define-range">Range Replay</a>. The CUDA context for which the range should be profiled must be current to the thread defining the start of the range and must be active for the entire range.</p>
<figure class="align-center" id="id17">
<img alt="../_images/replay-application-range.png" src="../_images/replay-application-range.png" />
<figcaption>
<p><span class="caption-text">Execution with Application Range Replay. An range of workloads is replayed by re-running the entire application without modifying interactions or saving and restoring memory.</span><a class="headerlink" href="#id17" title="Permalink to this image"></a></p>
</figcaption>
</figure>
</section>
<section id="graph-profiling">
<h4>Graph Profiling<a class="headerlink" href="#graph-profiling" title="Permalink to this headline"></a></h4>
<p>In multiple replay modes, NVIDIA Nsight Compute can profile CUDA graphs as single workload entities, rather than profile individual kernel nodes. The behavior can be toggled in the respective <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">command line</a> or <a class="reference external" href="../NsightCompute/index.html#connection-activity-interactive">UI</a> options.</p>
<p>The primary use cases for enabling this mode are:</p>
<ul class="simple">
<li><p>Profile graphs that include mandatory concurrent kernel nodes.</p></li>
<li><p>Profile graphs that include device-sided graph launches.</p></li>
<li><p>Profile graph behavior more accurately across multiple kernel node launches, as caches are not purged in between nodes.</p></li>
</ul>
<p>Note that when graph profiling is enabled, certain metrics such as instruction-level source metrics are not available. This then also applies to kernels profiled outside of graphs.</p>
</section>
</section>
<section id="compatibility">
<h3><span class="section-number">2.2.4. </span>Compatibility<a class="headerlink" href="#compatibility" title="Permalink to this headline"></a></h3>
<p>The set of available <a class="reference external" href="index.html#replay">replay modes</a> and <a class="reference external" href="index.html#metrics-guide">metrics</a> depends on the type of GPU workload to profile.</p>
<table class="table-no-stripes docutils align-default" id="id18">
<caption><span class="caption-text">Table 2. Replay modes and metric compatibility per workload type</span><a class="headerlink" href="#id18" title="Permalink to this table"></a></caption>
<colgroup>
<col style="width: 13%" />
<col style="width: 7%" />
<col style="width: 7%" />
<col style="width: 10%" />
<col style="width: 11%" />
<col style="width: 15%" />
<col style="width: 11%" />
<col style="width: 15%" />
<col style="width: 5%" />
<col style="width: 7%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p>Workload Type</p></td>
<td colspan="4"><p>Replay Mode</p></td>
<td colspan="5"><p>Metric Groups</p></td>
</tr>
<tr class="row-even"><td></td>
<td><p>Kernel</p></td>
<td><p>Application</p></td>
<td><p>Range</p></td>
<td><p>Application-Range</p></td>
<td><p>Hardware Counters / SMSP</p></td>
<td><p>Unit-Level Source</p></td>
<td><p>Instruction-Level Source</p></td>
<td><p>Launch</p></td>
<td><p>PM Sampling</p></td>
</tr>
<tr class="row-odd"><td><p>Kernel</p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>Yes <a class="footnote-reference brackets" href="#f2" id="id1">2</a></p></td>
<td><p>Yes <a class="footnote-reference brackets" href="#f2" id="id2">2</a></p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
</tr>
<tr class="row-even"><td><p>Range</p></td>
<td><p>No</p></td>
<td><p>No</p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>No</p></td>
<td><p>No</p></td>
<td><p>Some</p></td>
<td><p>Yes</p></td>
</tr>
<tr class="row-odd"><td><p>Cmdlist</p></td>
<td><p>Yes</p></td>
<td><p>No</p></td>
<td><p>No</p></td>
<td><p>Yes <a class="footnote-reference brackets" href="#f2" id="id3">2</a></p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>Yes</p></td>
<td><p>Some</p></td>
<td><p>Yes</p></td>
</tr>
<tr class="row-even"><td><p>Graph <a class="footnote-reference brackets" href="#f1" id="id4">1</a></p></td>
<td><p>Yes</p></td>
<td><p>No</p></td>
<td><p>No</p></td>
<td><p>Yes <a class="footnote-reference brackets" href="#f2" id="id5">2</a></p></td>
<td><p>Yes</p></td>
<td><p>No</p></td>
<td><p>No</p></td>
<td><p>Some</p></td>
<td><p>Yes</p></td>
</tr>
</tbody>
</table>
<p class="rubric">Footnotes</p>
<dl class="footnote brackets">
<dt class="label" id="f1"><span class="brackets"><a class="fn-backref" href="#id4">1</a></span></dt>
<dd><p>Limitations also apply to kernels profiled outside of graphs.</p>
</dd>
<dt class="label" id="f2"><span class="brackets">2</span><span class="fn-backref">(<a href="#id1">1</a>,<a href="#id2">2</a>,<a href="#id3">3</a>,<a href="#id5">4</a>)</span></dt>
<dd><p>Workload type is supported as part of the profiled range, but not separated in the result. Metric support matches that of Range workloads.</p>
</dd>
</dl>
</section>
<section id="profile-series">
<h3><span class="section-number">2.2.5. </span>Profile Series<a class="headerlink" href="#profile-series" title="Permalink to this headline"></a></h3>
<p>The performance of a kernel is highly dependent on the used launch parameters. Small changes to the launch parameters can have a significant effect on the runtime behavior of the kernel. However, identifying the best parameter set for a kernel by manually testing a lot of combinations can be a tedious process.</p>
<p>To make this workflow faster and more convenient, Profile Series provide the ability to automatically profile a single kernel multiple times with changing parameters. The parameters to be modified and values to be tested can be independently enabled and configured. For each combination of selected parameter values a unique profile result is collected. And the modified parameter values are tracked in the description of the results of a series. By comparing the results of a profile series, the kernel’s behavior on the changing parameters can be seen and the most optimal parameter set can be identified quickly.</p>
<figure class="align-center" id="id19">
<img alt="../_images/profile-series-action.png" src="../_images/profile-series-action.png" />
<figcaption>
<p><span class="caption-text">Profile Series action.</span><a class="headerlink" href="#id19" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<figure class="align-center" id="id20">
<img alt="../_images/profile-series-dialog.png" src="../_images/profile-series-dialog.png" />
<figcaption>
<p><span class="caption-text">Profile Series dialog.</span><a class="headerlink" href="#id20" title="Permalink to this image"></a></p>
</figcaption>
</figure>
</section>
<section id="overhead">
<h3><span class="section-number">2.2.6. </span>Overhead<a class="headerlink" href="#overhead" title="Permalink to this headline"></a></h3>
<p>As with most measurements, collecting performance data using NVIDIA Nsight Compute CLI incurs some runtime overhead on the application. The overhead does depend on a number of different factors:</p>
<ul>
<li><p><strong>Number and type of collected metrics</strong></p>
<p>Depending on the selected metric, data is collected either through a hardware performance monitor on the GPU, through software patching of the kernel instructions or via a launch or device attribute. The overhead between these mechanisms varies greatly, with launch and device attributes being “statically” available and requiring no kernel runtime overhead.</p>
<p>Furthermore, only a limited number of metrics can be collected in a single <em>pass</em> of the kernel execution. If more metrics are requested, the kernel launch is <em>replayed</em> multiple times, with its accessible memory being saved and restored between subsequent passes to guarantee deterministic execution. Therefore, collecting more metrics can significantly increase overhead by requiring more replay passes and increasing the total amount of memory that needs to be restored during replay.</p>
</li>
<li><p><strong>The collected section set</strong></p>
<p>Since each <a class="reference external" href="index.html#sets-and-sections">set</a> specifies a group of sections to be collected, choosing a less comprehensive set can reduce profiling overhead. See the <code class="docutils literal notranslate"><span class="pre">--set</span></code> command in the <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">NVIDIA Nsight Compute CLI</a> documentation.</p>
</li>
<li><p><strong>Number of collected sections</strong></p>
<p>Since each <a class="reference external" href="index.html#sets-and-sections">section</a> specifies a number of metrics to be collected, selecting fewer sections can reduce profiling overhead. See the <code class="docutils literal notranslate"><span class="pre">--section</span></code> command in the <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">NVIDIA Nsight Compute CLI</a> documentation.</p>
</li>
<li><p><strong>Number of profiled kernels</strong></p>
<p>By default, all selected metrics are collected for all launched kernels. To reduce the impact on the application, you can try to limit performance data collection to as few kernel functions and instances as makes sense for your analysis. See the filtering commands in the <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">NVIDIA Nsight Compute CLI</a> documentation.</p>
<p>There is a relatively high one-time overhead for the first profiled kernel in each context to generate the metric configuration. This overhead does not occur for subsequent kernels in the same context, if the list of collected metrics remains unchanged.</p>
</li>
<li><p><strong>GPU Architecture</strong></p>
<p>For some metrics, the overhead can vary depending on the exact chip they are collected on, e.g. due to varying number of units on the chip. Similarly, the overhead for resetting the L2 cache in-between kernel replay passes depends on the size of that cache.</p>
</li>
</ul>
</section>
</section>
<section id="metrics-guide">
<h2><span class="section-number">2.3. </span>Metrics Guide<a class="headerlink" href="#metrics-guide" title="Permalink to this headline"></a></h2>
<section id="hardware-model">
<span id="metrics-hw-model"></span><h3><span class="section-number">2.3.1. </span>Hardware Model<a class="headerlink" href="#hardware-model" title="Permalink to this headline"></a></h3>
<p class="rubric-h3 rubric">Compute Model</p>
<p>All NVIDIA GPUs are designed to support a general purpose heterogeneous parallel programming model, commonly known as <em>Compute</em>. This model decouples the GPU from the traditional graphics pipeline and exposes it as a general purpose parallel multi-processor. A heterogeneous computing model implies the existence of a host and a device, which in this case are the CPU and GPU, respectively. At a high level view, the host (CPU) manages resources between itself and the device and will send work off to the device to be executed in parallel.</p>
<p>Central to the compute model is the Grid, Block, Thread hierarchy, which defines how compute work is organized on the GPU. The hierarchy from top to bottom is as follows:</p>
<ul class="simple">
<li><p>A <em>Grid</em> is a 1D, 2D or 3D array of thread blocks.</p></li>
<li><p>A <em>Block</em> is a 1D, 2D or 3D array of threads, also known as a <em>Cooperative Thread Array (CTA)</em>.</p></li>
<li><p>A <em>Thread</em> is a single thread which runs on one of the GPU’s SM units.</p></li>
</ul>
<p>The purpose of the Grid, Block, Thread hierarchy is to expose a notion of locality amongst a group of threads, i.e. a Cooperative Thread Array (CTA). In CUDA, CTAs are referred to as Thread Blocks. The architecture can exploit this locality by providing fast shared memory and barriers between the threads within a single CTA. When a Grid is launched, the architecture guarantees that all threads within a CTA will run concurrently on the same SM. Information on the grids and blocks can be found in the <a class="reference external" href="index.html#sections-and-rules">Launch Statistics</a> section.</p>
<p>The number of CTAs that fit on each SM depends on the physical resources required by the CTA. These resource limiters include the number of threads and registers, shared memory utilization, and hardware barriers. The number CTAs per SM is referred to as the CTA <em>occupancy</em>, and these physical resources limit this occupancy. Details on the kernel’s occupancy are collected by the <a class="reference external" href="index.html#sections-and-rules">Occupancy</a> section.</p>
<p>Each CTA can be scheduled on any of the available SMs, where there is no guarantee in the order of execution. As such, CTAs must be entirely independent, which means it is not possible for one CTA to wait on the result of another CTA. As CTAs are independent, the host (CPU) can launch a large Grid that will not fit on the hardware all at once, however any GPU will still be able to run it and produce the correct results.</p>
<p>CTAs are further divided into groups of 32 threads called <em>Warps</em>. If the number of threads in a CTA is not dividable by 32, the last warp will contain the remaining number of threads.</p>
<p>The total number of CTAs that can run concurrently on a given GPU is referred to as <em>Wave</em>. Consequently, the size of a Wave scales with the number of available SMs of a GPU, but also with the occupancy of the kernel.</p>
<p class="rubric-h3 rubric">Streaming Multiprocessor</p>
<p>The <em>Streaming Multiprocessor (SM)</em> is the core processing unit in the GPU. The SM is optimized for a wide diversity of workloads, including general-purpose computations, deep learning, ray tracing, as well as lighting and shading. The SM is designed to simultaneously execute multiple CTAs. CTAs can be from different grid launches.</p>
<p>The SM implements an execution model called Single Instruction Multiple Threads (SIMT), which allows individual threads to have unique control flow while still executing as part of a warp. The Turing SM inherits the Volta SM’s independent thread scheduling model. The SM maintains execution state per thread, including a program counter (PC) and call stack. The independent thread scheduling allows the GPU to yield execution of any thread, either to make better use of execution resources or to allow a thread to wait for data produced by another thread possibly in the same warp. Collecting the <a class="reference external" href="index.html#sections-and-rules">Source Counters</a> section allows you to inspect instruction execution and predication details on the <em>Source Page</em>, along with <a class="reference external" href="index.html#sampling">Sampling</a> information.</p>
<p>Each SM is partitioned into four processing blocks, called <em>SM sub partitions</em>. The SM sub partitions are the primary processing elements on the SM. Each sub partition contains the following units:</p>
<ul class="simple">
<li><p>Warp Scheduler</p></li>
<li><p>Register File</p></li>
<li><p>Execution Units/Pipelines/Cores</p>
<ul>
<li><p>Integer Execution units</p></li>
<li><p>Floating Point Execution units</p></li>
<li><p>Memory Load/Store units</p></li>
<li><p>Special Function unit</p></li>
<li><p>Tensor Cores</p></li>
</ul>
</li>
</ul>
<p>Shared within an SM across the four SM partitions are:</p>
<ul class="simple">
<li><p>Unified L1 Data Cache / Shared Memory</p></li>
<li><p>Texture units</p></li>
<li><p>RT Cores, if available</p></li>
</ul>
<p>A warp is allocated to a sub partition and resides on the sub partition from launch to completion. A warp is referred to as <em>active</em> or <em>resident</em> when it is mapped to a sub partition. A sub partition manages a fixed size pool of warps. On Volta architectures, the size of the pool is 16 warps. On Turing architectures the size of the pool is 8 warps. Active warps can be in <em>eligible</em> state if the warp is ready to issue an instruction. This requires the warp to have a decoded instruction, all input dependencies resolved, and for the function unit to be available. Statistics on active, eligible and issuing warps can be collected with the <a class="reference external" href="index.html#sections-and-rules">Scheduler Statistics</a> section.</p>
<p>A warp is <em>stalled</em> when the warp is waiting on</p>
<ul class="simple">
<li><p>an instruction fetch,</p></li>
<li><p>a memory dependency (result of memory instruction),</p></li>
<li><p>an execution dependency (result of previous instruction), or</p></li>
<li><p>a synchronization barrier.</p></li>
</ul>
<p>See <a class="reference external" href="index.html#statistical-sampler">Warp Scheduler States</a> for the list of stall reasons that can be profiled and the <a class="reference external" href="index.html#sections-and-rules">Warp State Statistics</a> section for a summary of warp states found in the kernel execution.</p>
<p>The most important resource under the compiler’s control is the number of registers used by a kernel. Each sub partition has a set of 32-bit registers, which are allocated by the HW in fixed-size chunks. The <a class="reference external" href="index.html#sections-and-rules">Launch Statistics</a> section shows the kernel’s register usage.</p>
<p class="rubric-h3 rubric" id="metrics-hw-memory">Memory</p>
<p>Global memory is a 49-bit virtual address space that is mapped to physical memory on the device, pinned system memory, or peer memory. Global memory is visible to all threads in the GPU. Global memory is accessed through the SM L1 and GPU L2.</p>
<p>Local memory is private storage for an executing thread and is not visible outside of that thread. It is intended for thread-local data like thread stacks and register spills. Local memory addresses are translated to global virtual addresses by the the AGU unit. Local memory has the same latency as global memory. One difference between global and local memory is that local memory is arranged such that consecutive 32-bit words are accessed by consecutive thread IDs. Accesses are therefore fully coalesced as long as all threads in a warp access the same relative address (e.g., same index in an array variable, same member in a structure variable, etc.).</p>
<p>Shared memory is located on chip, so it has much higher bandwidth and much lower latency than either local or global memory. Shared memory can be shared across a compute CTA. Compute CTAs attempting to share data across threads via shared memory must use synchronization operations (such as __syncthreads()) between stores and loads to ensure data written by any one thread is visible to other threads in the CTA. Similarly, threads that need to share data via global memory must use a more heavyweight global memory barrier.</p>
<p>Shared memory has 32 banks that are organized such that successive 32-bit words map to successive banks that can be accessed simultaneously. Any 32-bit memory read or write request made of 32 addresses that fall in 32 distinct memory banks can therefore be serviced simultaneously, yielding an overall bandwidth that is 32 times as high as the bandwidth of a single request. However, if two addresses of a memory request fall in the same memory bank, there is a bank conflict and the access has to be serialized.</p>
<p>A shared memory request for a warp does not generate a bank conflict between two threads that access any address within the same 32-bit word (even though the two addresses fall in the same bank). When multiple threads make the same read access, one thread receives the data and then broadcasts it to the other threads. When multiple threads write to the same location, only one thread succeeds in the write; which thread that succeeds is undefined.</p>
<p>Detailed memory metrics are collected by the <a class="reference external" href="index.html#sections-and-rules">Memory Workload Analysis</a> section.</p>
<p class="rubric-h3 rubric" id="metrics-hw-caches">Caches</p>
<p>All GPU units communicate to main memory through the Level 2 cache, also known as the L2. The L2 cache sits between on-chip memory clients and the framebuffer. L2 works in physical-address space. In addition to providing caching functionality, L2 also includes hardware to perform compression and global atomics.</p>
<figure class="align-center" id="id21">
<img alt="../_images/hw-model-lts.png" src="../_images/hw-model-lts.png" />
<figcaption>
<p><span class="caption-text">Model of the L2 cache.</span><a class="headerlink" href="#id21" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p>The Level 1 Data Cache, or L1, plays a key role in handling global, local, shared, texture, and surface memory reads and writes, as well as reduction and atomic operations. On Volta and Turing architectures there are , there are two L1 caches per TPC, one for each SM. For more information on how L1 fits into the texturing pipeline, see the <a class="reference external" href="index.html#metrics-hw-tex-surf">TEX unit</a> description. Also note that while this section often uses the name “L1”, it should be understood that the L1 data cache, shared data, and the Texture data cache are one and the same.</p>
<p>L1 receives requests from two units: the SM and TEX. L1 receives global and local memory requests from the SM and receives texture and surface requests from TEX. These operations access memory in the global memory space, which L1 sends through a secondary cache, the L2.</p>
<p>Cache hit and miss rates as well as data transfers are reported in the <a class="reference external" href="index.html#sections-and-rules">Memory Workload Analysis</a> section.</p>
<figure class="align-center" id="id22">
<img alt="../_images/hw-model-l1tex.png" src="../_images/hw-model-l1tex.png" />
<figcaption>
<p><span class="caption-text">Model of Load/Store and Texture pipelines for the L1TEX cache.</span><a class="headerlink" href="#id22" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="rubric-h3 rubric" id="metrics-hw-tex-surf">Texture/Surface</p>
<p>The TEX unit performs texture fetching and filtering. Beyond plain texture memory access, TEX is responsible for the addressing, LOD, wrap, filter, and format conversion operations necessary to convert a texture read request into a result.</p>
<p>TEX receives two general categories of requests from the SM via its input interface: texture requests and surface load/store operations. Texture and surface memory space resides in device memory and are cached in L1. Texture and surface memory are allocated as block-linear surfaces (e.g. 2D, 2D Array, 3D). Such surfaces provide a cache-friendly layout of data such that neighboring points on a 2D surface are also located close to each other in memory, which improves access locality. Surface accesses are bounds-checked by the TEX unit prior to accessing memory, which can be used for implementing different texture wrapping modes.</p>
<p>The L1 cache is optimized for 2D spatial locality, so threads of the same warp that read texture or surface addresses that are close together in 2D space will achieve optimal performance. The L1 cache is also designed for streaming fetches with constant latency; a cache hit reduces DRAM bandwidth demand but not fetch latency. Reading device memory through texture or surface memory presents some benefits that can make it an advantageous alternative to reading memory from global or constant memory.</p>
<p>Information on texture and surface memory can be found in the <a class="reference external" href="index.html#sections-and-rules">Memory Workload Analysis</a> section.</p>
</section>
<section id="metrics-structure">
<h3><span class="section-number">2.3.2. </span>Metrics Structure<a class="headerlink" href="#metrics-structure" title="Permalink to this headline"></a></h3>
<p class="rubric-h3 rubric" id="metrics-overview">Metrics Overview</p>
<p>NVIDIA Nsight Compute uses an advanced metrics calculation system, designed to help you determine what happened (counters and metrics), and how close the program reached to peak GPU performance (throughputs as a percentage). Every counter has associated peak rates in the database, to allow computing its throughput as a percentage.</p>
<p>Throughput metrics return the maximum percentage value of their constituent counters. These constituents have been carefully selected to represent the sections of the GPU pipeline that govern peak performance. While all counters can be converted to a %-of-peak, not all counters are suitable for peak-performance analysis; examples of unsuitable counters include qualified subsets of activity, and workload residency counters. Using throughput metrics ensures meaningful and actionable analysis.</p>
<p>Two types of peak rates are available for every counter: burst and sustained. Burst rate is the maximum rate reportable in a single clock cycle. Sustained rate is the maximum rate achievable over an infinitely long measurement period, for “typical” operations. For many counters, burst equals sustained. Since the burst rate cannot be exceeded, percentages of burst rate will always be less than 100%. Percentages of sustained rate can occasionally exceed 100% in edge cases.</p>
<p class="rubric-h3 rubric" id="metrics-entities">Metrics Entities</p>
<p>While in NVIDIA Nsight Compute, all performance counters are named <em>metrics</em>, they can be split further into groups with specific properties. For metrics collected via the <em>PerfWorks</em> measurement library, the following entities exist:</p>
<p><strong>Counters</strong> may be either a raw counter from the GPU, or a calculated counter value. Every counter has four sub-metrics under it, which are also called <em>roll-ups</em>:</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 13%" />
<col style="width: 87%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.sum</span></code></p></td>
<td><p>The sum of counter values across all unit instances.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.avg</span></code></p></td>
<td><p>The average counter value across all unit instances.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.min</span></code></p></td>
<td><p>The minimum counter value across all unit instances.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.max</span></code></p></td>
<td><p>The maximum counter value across all unit instances.</p></td>
</tr>
</tbody>
</table>
<p>Counter roll-ups have the following calculated quantities as built-in sub-metrics:</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 36%" />
<col style="width: 64%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained</span></code></p></td>
<td><p>the peak sustained rate</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_active</span></code></p></td>
<td><p>the peak sustained rate during unit active cycles</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_active.per_second</span></code></p></td>
<td><p>the peak sustained rate during unit active cycles, per second *</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_elapsed</span></code></p></td>
<td><p>the peak sustained rate during unit elapsed cycles</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_elapsed.per_second</span></code></p></td>
<td><p>the peak sustained rate during unit elapsed cycles, per second *</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.per_second</span></code></p></td>
<td><p>the number of operations per second</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.per_cycle_active</span></code></p></td>
<td><p>the number of operations per unit active cycle</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.per_cycle_elapsed</span></code></p></td>
<td><p>the number of operations per unit elapsed cycle</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_active</span></code></p></td>
<td><p>% of peak sustained rate achieved during unit active cycles</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_elapsed</span></code></p></td>
<td><p>% of peak sustained rate achieved during unit elapsed cycles</p></td>
</tr>
</tbody>
</table>
<p>* sub-metrics added in NVIDIA Nsight Compute 2022.2.0.</p>
<p>Example: <code class="docutils literal notranslate"><span class="pre">ncu</span> <span class="pre">--query-metrics-mode</span> <span class="pre">suffix</span> <span class="pre">--metrics</span> <span class="pre">sm__inst_executed</span> <span class="pre">--chip</span> <span class="pre">ga100</span></code></p>
<p><strong>Ratios</strong> have three sub-metrics:</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 27%" />
<col style="width: 73%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct</span></code></p></td>
<td><p>The value expressed as a percentage.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.ratio</span></code></p></td>
<td><p>The value expressed as a ratio.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.max_rate</span></code></p></td>
<td><p>The ratio’s maximum value.</p></td>
</tr>
</tbody>
</table>
<p>Example: <code class="docutils literal notranslate"><span class="pre">ncu</span> <span class="pre">--query-metrics-mode</span> <span class="pre">suffix</span> <span class="pre">--metrics</span> <span class="pre">smsp__average_warp_latency</span> <span class="pre">--chip</span> <span class="pre">ga100</span></code></p>
<p><strong>Throughputs</strong> indicate how close a portion of the GPU reached to peak rate. Every throughput has the following sub-metrics:</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 34%" />
<col style="width: 66%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_active</span></code></p></td>
<td><p>% of peak sustained rate achieved during unit active cycles</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_elapsed</span></code></p></td>
<td><p>% of peak sustained rate achieved during unit elapsed cycles</p></td>
</tr>
</tbody>
</table>
<p>Example: <code class="docutils literal notranslate"><span class="pre">ncu</span> <span class="pre">--query-metrics-mode</span> <span class="pre">suffix</span> <span class="pre">--metrics</span> <span class="pre">sm__throughput</span> <span class="pre">--chip</span> <span class="pre">ga100</span></code></p>
<p>Throughputs have a breakdown of underlying metrics from which the throughput value is computed. You can collect <code class="docutils literal notranslate"><span class="pre">breakdown:<throughput-metric></span></code> to collect a throughput’s breakdown metrics.</p>
<p><strong>Deprecated counter sub-metrics:</strong> The following sub-metrics were removed, due to not being useful for performance optimization:</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 36%" />
<col style="width: 64%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_burst</span></code></p></td>
<td><p>the peak burst rate</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_active</span></code></p></td>
<td><p>% of peak burst rate achieved during unit active cycles</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_elapsed</span></code></p></td>
<td><p>% of peak burst rate achieved during unit elapsed cycles</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_region</span></code></p></td>
<td><p>% of peak burst rate achieved over a user-specified “range”</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_frame</span></code></p></td>
<td><p>% of peak burst rate achieved over a user-specified “frame”</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_region</span></code></p></td>
<td><p>% of peak sustained rate achieved over a user-specified “range” time</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_frame</span></code></p></td>
<td><p>% of peak sustained rate achieved over a user-specified “frame” time</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.per_cycle_in_region</span></code></p></td>
<td><p>the number of operations per user-specified “range” cycle</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.per_cycle_in_frame</span></code></p></td>
<td><p>the number of operations per user-specified “frame” cycle</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_region</span></code></p></td>
<td><p>the peak sustained rate over a user-specified “range”</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_region.per_second</span></code></p></td>
<td><p>the peak sustained rate over a user-specified “range”, per second *</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_frame</span></code></p></td>
<td><p>the peak sustained rate over a user-specified “frame”</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.peak_sustained_frame.per_second</span></code></p></td>
<td><p>the peak sustained rate over a user-specified “frame”, per second *</p></td>
</tr>
</tbody>
</table>
<p><strong>Deprecated throughput sub-metrics:</strong> The following sub-metrics were removed, due to not being useful for performance optimization:</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 36%" />
<col style="width: 64%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_active</span></code></p></td>
<td><p>% of peak burst rate achieved during unit active cycles</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_elapsed</span></code></p></td>
<td><p>% of peak burst rate achieved during unit elapsed cycles</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_region</span></code></p></td>
<td><p>% of peak burst rate achieved over a user-specified “range” time</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_burst_frame</span></code></p></td>
<td><p>% of peak burst rate achieved over a user-specified “frame” time</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_region</span></code></p></td>
<td><p>% of peak sustained rate achieved over a user-specified “range”</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">.pct_of_peak_sustained_frame</span></code></p></td>
<td><p>% of peak sustained rate achieved over a user-specified “frame”</p></td>
</tr>
</tbody>
</table>
<p>In addition to PerfWorks metrics, NVIDIA Nsight Compute uses several other measurement providers that each generate their own metrics. These are explained in the <a class="reference external" href="index.html#metrics-reference">Metrics Reference</a>.</p>
<p class="rubric-h3 rubric" id="metrics-examples">Metrics Examples</p>
<div class="highlight-text notranslate"><div class="highlight"><pre><span></span>## non-metric names -- *not* directly evaluable
sm__inst_executed # counter
smsp__average_warp_latency # ratio
sm__throughput # throughput
## a counter's four first-level sub-metrics -- all evaluable
sm__inst_executed.sum
sm__inst_executed.avg
sm__inst_executed.min
sm__inst_executed.max
## all names below are metrics -- all evaluable
l1tex__data_bank_conflicts_pipe_lsu.sum
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained_active
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained_active.per_second
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained_elapsed
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained_elapsed.per_second
l1tex__data_bank_conflicts_pipe_lsu.sum.per_cycle_active
l1tex__data_bank_conflicts_pipe_lsu.sum.per_cycle_elapsed
l1tex__data_bank_conflicts_pipe_lsu.sum.per_second
l1tex__data_bank_conflicts_pipe_lsu.sum.pct_of_peak_sustained_active
l1tex__data_bank_conflicts_pipe_lsu.sum.pct_of_peak_sustained_elapsed
...
</pre></div>
</div>
<p class="rubric-h3 rubric" id="metrics-naming-conventions">Metrics Naming Conventions</p>
<p>Counters and metrics _generally_ obey the naming scheme:</p>
<ul class="simple">
<li><p>Unit-Level Counter : <code class="docutils literal notranslate"><span class="pre">unit__(subunit?)_(pipestage?)_quantity_(qualifiers?)</span></code></p></li>
<li><p>Interface Counter : <code class="docutils literal notranslate"><span class="pre">unit__(subunit?)_(pipestage?)_(interface)_quantity_(qualifiers?)</span></code></p></li>
<li><p>Unit Metric : <code class="docutils literal notranslate"><span class="pre">(counter_name).(rollup_metric)</span></code></p></li>
<li><p>Sub-Metric : <code class="docutils literal notranslate"><span class="pre">(counter_name).(rollup_metric).(submetric)</span></code></p></li>
</ul>
<p>where</p>
<ul class="simple">
<li><p>unit: A logical or physical unit of the GPU</p></li>
<li><p>subunit: The subunit within the unit where the counter was measured. Sometimes this is a pipeline mode instead.</p></li>
<li><p>pipestage: The pipeline stage within the subunit where the counter was measured.</p></li>
<li><p>quantity: What is being measured. Generally matches the <em>dimensional units</em>.</p></li>
<li><p>qualifiers: Any additional predicates or filters applied to the counter. Often, an unqualified counter can be broken down into several qualified sub-components.</p></li>
<li><p>interface: Of the form <code class="docutils literal notranslate"><span class="pre">sender2receiver</span></code>, where <code class="docutils literal notranslate"><span class="pre">sender</span></code> is the source-unit and <code class="docutils literal notranslate"><span class="pre">receiver</span></code> is the destination-unit.</p></li>
<li><p>rollup_metric: One of sum, avg, min, max.</p></li>
<li><p>submetric: refer to section <a class="reference external" href="index.html#metrics-entities">Metrics Entities</a></p></li>
</ul>
<p>Components are not always present. Most top-level counters have no qualifiers. Subunit and pipestage may be absent where irrelevant, or there may be many subunit specifiers for detailed counters.</p>
<p class="rubric-h3 rubric" id="cycle-metrics">Cycle Metrics</p>
<p>Counters using the term <code class="docutils literal notranslate"><span class="pre">cycles</span></code> in the name report the number of cycles in the unit’s clock domain. Unit-level cycle metrics include:</p>
<ul class="simple">
<li><p><code class="docutils literal notranslate"><span class="pre">unit__cycles_elapsed</span></code> : The number of cycles within a range. The cycles’ DimUnits are specific to the unit’s clock domain.</p></li>
<li><p><code class="docutils literal notranslate"><span class="pre">unit__cycles_active</span></code> : The number of cycles where the unit was processing data.</p></li>
<li><p><code class="docutils literal notranslate"><span class="pre">unit__cycles_stalled</span></code> : The number of cycles where the unit was unable to process new data because its output interface was blocked.</p></li>
<li><p><code class="docutils literal notranslate"><span class="pre">unit__cycles_idle</span></code> : The number of cycles where the unit was idle.</p></li>
</ul>
<p>Interface-level cycle counters are often (not always) available in the following variations:</p>
<ul class="simple">
<li><p><code class="docutils literal notranslate"><span class="pre">unit__(interface)_active</span></code> : Cycles where data was transferred from source-unit to destination-unit.</p></li>
<li><p><code class="docutils literal notranslate"><span class="pre">unit__(interface)_stalled</span></code> : Cycles where the source-unit had data, but the destination-unit was unable to accept data.</p></li>
</ul>
<p class="rubric-h3 rubric" id="instanced-metrics">Instanced Metrics</p>
<p>Metrics collected with NVIDIA Nsight Compute can have a single (aggregate) value, multiple instance values, or both. Instances allow the metric to have multiple sub-values, e.g. representing the value of an source metric at each instruction offset. If a metric has instance values, it often also has a correlation ID for each instance. Correlation IDs and values form a mapping that allows the tool to correlate the values within a context. For source metrics, that context is commonly the address ranges of the functions executed as part of the workload.</p>
<p>You can find which metrics have instance values in the <a class="reference external" href="index.html#metrics-reference__metrics-ref-source">Metrics Reference</a>. In the UI, the <a class="reference external" href="../NsightCompute/index.html#tool-window-metric-details">Metric Details</a> tool window can be used to conveniently view correlation IDs and instance values for each metric. Also, both the UI and the command line interface provide options to show instance values in addition to a metric aggregate where applicable.</p>
</section>
<section id="metrics-decoder">
<h3><span class="section-number">2.3.3. </span>Metrics Decoder<a class="headerlink" href="#metrics-decoder" title="Permalink to this headline"></a></h3>
<p>The following explains terms found in NVIDIA Nsight Compute metric names, as introduced in <a class="reference external" href="index.html#metrics-structure">Metrics Structure</a>.</p>
<blockquote>
<div><p class="rubric-h3 rubric" id="metrics-hw-units">Units</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id23">
<caption><span class="caption-text">Units</span><a class="headerlink" href="#id23" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">dram</span></code></p></td>
<td><p>Device (main) memory, where the GPUs global and local memory resides.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">fbpa</span></code></p></td>
<td><p>The FrameBuffer Partition is a memory controller which sits between the level 2 cache (LTC) and the DRAM. The number of FBPAs varies across GPUs.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">fe</span></code></p></td>
<td><p>The Frontend unit is responsible for the overall flow of workloads sent by the driver. FE also facilitates a number of synchronization operations.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">gpc</span></code></p></td>
<td><p>The General Processing Cluster contains SM, Texture and L1 in the form of TPC(s). It is replicated several times across a chip.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">gpu</span></code></p></td>
<td><p>The entire Graphics Processing Unit.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">gr</span></code></p></td>
<td><p>Graphics Engine is responsible for all 2D and 3D graphics, compute work, and synchronous graphics copying work.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">idc</span></code></p></td>
<td><p>The InDexed Constant Cache is a subunit of the SM responsible for caching constants that are indexed with a register.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">l1tex</span></code></p></td>
<td><p>The Level 1 (L1)/Texture Cache is located within the GPC. It can be used as directed-mapped shared memory and/or store global, local and texture data in its cache portion. l1tex__t refers to its Tag stage. l1tex__m refers to its Miss stage. l1tex__d refers to its Data stage.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">ltc</span></code></p></td>
<td><p>The Level 2 cache.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">ltcfabric</span></code></p></td>
<td><p>The LTC fabric is the communication fabric for the L2 cache partitions.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">lts</span></code></p></td>
<td><p>A Level 2 (L2) Cache Slice is a sub-partition of the Level 2 cache. lts__t refers to its Tag stage. lts__m refers to its Miss stage. lts__d refers to its Data stage.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">mcc</span></code></p></td>
<td><p>Memory controller channel of MSS. The Memory Subsystem (MSS) provides access to local DRAM, SysRAM, and provides a SyncPoint Interface for interprocessor signaling. MCC includes the row sorter/arbiter and DRAM controllers.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">pm</span></code></p></td>
<td><p>Performance monitor.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">sm</span></code></p></td>
<td><p>The Streaming Multiprocessor handles execution of a kernel as groups of 32 threads, called warps. Warps are further grouped into cooperative thread arrays (CTA), called blocks in CUDA. All warps of a CTA execute on the same SM. CTAs share various resources across their threads, e.g. the shared memory.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp</span></code></p></td>
<td><p>Each SM is partitioned into four processing blocks, called SM sub partitions. The SM sub partitions are the primary processing elements on the SM. A sub partition manages a fixed size pool of warps.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">sys</span></code></p></td>
<td><p>Logical grouping of several units.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">tpc</span></code></p></td>
<td><p>Thread Processing Clusters are units in the GPC. They contain one or more SM, Texture and L1 units, the Instruction Cache (ICC) and the Indexed Constant Cache (IDC).</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-sub-units">Subunits</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id24">
<caption><span class="caption-text">Subunits</span><a class="headerlink" href="#id24" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">aperture_device</span></code></p></td>
<td><p>Memory interface to local device memory (dram)</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">aperture_peer</span></code></p></td>
<td><p>Memory interface to remote device memory</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">aperture_sysmem</span></code></p></td>
<td><p>Memory interface to system memory</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">global</span></code></p></td>
<td><p>Global memory is a 49-bit virtual address space that is mapped to physical memory on the device, pinned system memory, or peer memory. Global memory is visible to all threads in the GPU. Global memory is accessed through the SM L1 and GPU L2.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">lg</span></code></p></td>
<td><p>Local/Global memory</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">local</span></code></p></td>
<td><p>Local memory is private storage for an executing thread and is not visible outside of that thread. It is intended for thread-local data like thread stacks and register spills. Local memory has the same latency as global memory.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">lsu</span></code></p></td>
<td><p>Load/Store unit</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">lsuin</span></code></p></td>
<td><p>Load/Store input</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">mio</span></code></p></td>
<td><p>Memory input/output</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">mioc</span></code></p></td>
<td><p>Memory input/output control</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">shared</span></code></p></td>
<td><p>Shared memory is located on chip, so it has much higher bandwidth and much lower latency than either local or global memory. Shared memory can be shared across a compute CTA.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">surface</span></code></p></td>
<td><p>Surface memory</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">texin</span></code></p></td>
<td><p>TEXIN</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">texture</span></code></p></td>
<td><p>Texture memory</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">xbar</span></code></p></td>
<td><p>The Crossbar (XBAR) is responsible for carrying packets from a given source unit to a specific destination unit.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-hw-pipelines">Pipelines</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id25">
<caption><span class="caption-text">Pipelines</span><a class="headerlink" href="#id25" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">adu</span></code></p></td>
<td><p>Address Divergence Unit. The ADU is responsible for address divergence handling for branches/jumps. It also provides support for constant loads and block-level barrier instructions.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">alu</span></code></p></td>
<td><p>Arithmetic Logic Unit. The ALU is responsible for execution of most bit manipulation and logic instructions. It also executes integer instructions, excluding IMAD and IMUL. On NVIDIA Ampere architecture chips, the ALU pipeline performs fast FP32-to-FP16 conversion.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">cbu</span></code></p></td>
<td><p>Convergence Barrier Unit. The CBU is responsible for warp-level convergence, barrier, and branch instructions.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">fma</span></code></p></td>
<td><p>Fused Multiply Add/Accumulate. The FMA pipeline processes most FP32 arithmetic (FADD, FMUL, FMAD). It also performs integer multiplication operations (IMUL, IMAD), as well as integer dot products. On GA10x, FMA is a logical pipeline that indicates peak FP32 and FP16x2 performance. It is composed of the FMAHeavy and FMALite physical pipelines.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">fmaheavy</span></code></p></td>
<td><p>Fused Multiply Add/Accumulate Heavy. FMAHeavy performs FP32 arithmetic (FADD, FMUL, FMAD), FP16 arithmetic (HADD2, HMUL2, HFMA2), integer multiplication operations (IMUL, IMAD), and integer dot products.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">fmalite</span></code></p></td>
<td><p>Fused Multiply Add/Accumulate Lite. FMALite performs FP32 arithmetic (FADD, FMUL, FMA) and FP16 arithmetic (HADD2, HMUL2, HFMA2).</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">fp16</span></code></p></td>
<td><p>Half-precision floating-point. On Volta, Turing and NVIDIA GA100, the FP16 pipeline performs paired FP16 instructions (FP16x2). It also contains a fast FP32-to-FP16 and FP16-to-FP32 converter. Starting with GA10x chips, this functionality is part of the FMA pipeline.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">fp64</span></code></p></td>
<td><p>Double-precision floating-point. The implementation of FP64 varies greatly per chip.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">lsu</span></code></p></td>
<td><p>Load Store Unit. The LSU pipeline issues load, store, atomic, and reduction instructions to the L1TEX unit for global, local, and shared memory. It also issues special register reads (S2R), shuffles, and CTA-level arrive/wait barrier instructions to the L1TEX unit.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">tex</span></code></p></td>
<td><p>Texture Unit. The SM texture pipeline forwards texture and surface instructions to the L1TEX unit’s TEXIN stage. On GPUs where FP64 or Tensor pipelines are decoupled, the texture pipeline forwards those types of instructions, too.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">tma</span></code></p></td>
<td><p>Tensor Memory Access Unit. Provides efficient data transfer mechanisms between global and shared memories with the ability to understand and traverse multidimensional data layouts.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">uniform</span></code></p></td>
<td><p>Uniform Data Path. This scalar unit executes instructions where all threads use the same input and generate the same output.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">xu</span></code></p></td>
<td><p>Transcendental and Data Type Conversion Unit. The XU pipeline is responsible for special functions such as sin, cos, and reciprocal square root. It is also responsible for int-to-float, and float-to-int type conversions.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-quantities">Quantities</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id26">
<caption><span class="caption-text">Quantities</span><a class="headerlink" href="#id26" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">instruction</span></code></p></td>
<td><p>An assembly (SASS) instruction. Each executed instruction may generate zero or more requests.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">request</span></code></p></td>
<td><p>A command into a HW unit to perform some action, e.g. load data from some memory location. Each request accesses one or more sectors.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">sector</span></code></p></td>
<td><p>Aligned 32 byte-chunk of memory in a cache line or device memory. An L1 or L2 cache line is four sectors, i.e. 128 bytes. Sector accesses are classified as hits if the tag is present and the sector-data is present within the cache line. Tag-misses and tag-hit-data-misses are all classified as misses.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">tag</span></code></p></td>
<td><p>Unique key to a cache line. A request may look up multiple tags, if the thread addresses do not all fall within a single cache line-aligned region. The L1 and L2 both have 128 byte cache lines. Tag accesses may be classified as hits or misses.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">wavefront</span></code></p></td>
<td><p>Unique “work package” generated at the end of the processing stage for requests. All work items of a wavefront are processed in parallel, while work items of different wavefronts are serialized and processed on different cycles. At least one wavefront is generated for each request.</p></td>
</tr>
</tbody>
</table>
</div></blockquote>
<p>A simplified model for the processing in L1TEX for Volta and newer architectures can be described as follows: When an SM executes a global or local memory instruction for a warp, a single <em>request</em> is sent to L1TEX. This request communicates the information for all participating threads of this warp (up to 32). For local and global memory, based on the access pattern and the participating threads, the request requires to access a number of cache lines, and <em>sectors</em> within these cache lines. The L1TEX unit has internally multiple processing stages operating in a pipeline.</p>
<p>A <em>wavefront</em> is the maximum unit that can pass through that pipeline stage per cycle. If not all cache lines or sectors can be accessed in a single wavefront, multiple wavefronts are created and sent for processing one by one, i.e. in a serialized manner. Limitations of the work within a wavefront may include the need for a consistent memory space, a maximum number of cache lines that can be accessed, as well as various other reasons. Each wavefront then flows through the L1TEX pipeline and fetches the sectors handled in that wavefront. The given relationships of the three key values in this model are <em>requests:sectors is 1:N, wavefronts:sectors 1:N, and requests:wavefronts is 1:N</em>.</p>
<p>A wavefront is described as a (work) package that can be processed at once, i.e. there is a notion of processing one wavefront per cycle in L1TEX. Wavefronts therefore represent the number of cycles required to process the requests, while the number of sectors per request is a property of the <em>access pattern</em> of the memory instruction for all participating threads. For example, it is possible to have a memory instruction that requires 4 sectors per request in 1 wavefront. However, you can also have a memory instruction having 4 sectors per request, but requiring 2 or more wavefronts.</p>
</section>
<section id="range-and-precision">
<h3><span class="section-number">2.3.4. </span>Range and Precision<a class="headerlink" href="#range-and-precision" title="Permalink to this headline"></a></h3>
<p class="title sectiontitle rubric" id="overview">Overview</p>
<p>In general, measurement values that lie outside the expected logical range of a metric can be attributed to one or more of the below root-causes. If values are exceeding such range, they are not clamped by the tool to their expected value on purpose to ensure that the rest of the profiler report remains self-consistent.</p>
<p class="title sectiontitle rubric" id="asynchronous-gpu-activity">Asynchronous GPU activity</p>
<p>GPU engines other than the one measured by a metric (display, copy engine, video encoder, video decoder, etc.) potentially access shared resources during profiling. Such chip-global shared resources include L2, DRAM, PCIe, and NVLINK. If the kernel launch is small, the other engine(s) can cause significant confusion in e.g. the DRAM results, since it is not possible to isolate the DRAM traffic of the SM. To reduce the impact of such asynchronous units, consider profiling on a GPU without active display and without other processes that can access the GPU at the time.</p>
<p class="title sectiontitle rubric" id="multi-pass-data-collection">Multi-pass data collection</p>
<p>Out-of-range metrics often occur when the profiler <a class="reference external" href="index.html#kernel-replay">replays</a> the kernel launch to collect metrics, and work distribution is significantly different across replay passes. A metric such as hit rate (hits / queries) can have significant error if hits and queries are collected on different passes and the kernel does not saturate the GPU to reach a steady state (generally > 20 µs). Similarly, it can show unexpected values when the workload is inherently variable, as e.g. in the case of spin loops.</p>
<p>To mitigate the issue, when applicable try to increase the measured workload to allow the GPU to reach a steady state for each launch. Reducing the number of metrics collected at the same time can also improve precision by increasing the likelihood that counters contributing to one metric are collected in a single pass.</p>
<p class="title sectiontitle rubric" id="tool-issue">Tool issue</p>
<p>If you still observe metric issues after following the guidelines above, please <a class="reference external" href="https://forums.developer.nvidia.com/c/development-tools/nsight-compute">reach out to us</a> and describe your issue.</p>
</section>
</section>
<section id="metrics-reference">
<h2><span class="section-number">2.4. </span>Metrics Reference<a class="headerlink" href="#metrics-reference" title="Permalink to this headline"></a></h2>
<p class="rubric-h2 rubric" id="overview-1">Overview</p>
<p>Most metrics in NVIDIA Nsight Compute can be queried using the ncu command line interface’s <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">–query-metrics</a> option.</p>
<p>The following metrics can be collected explicitly, but are not listed by <code class="docutils literal notranslate"><span class="pre">--query-metrics</span></code>, and do not follow the naming scheme explained in <a class="reference external" href="index.html#metrics-structure">Metrics Structure</a>. They should be used as-is instead.</p>
<p><code class="docutils literal notranslate"><span class="pre">launch__*</span></code> metrics are collected per kernel launch, and do not require an additional replay pass. They are available as part of the kernel launch parameters (such as grid size, block size, …) or are computed using the <a class="reference external" href="../NsightCompute/index.html#occupancy-calculator">CUDA Occupancy Calculator</a>.</p>
<blockquote>
<div><p class="rubric-h3 rubric" id="metrics-ref-launch">Launch Metrics</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id27">
<caption><span class="caption-text">Launch Metrics</span><a class="headerlink" href="#id27" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__block_dim_x</span></code></p></td>
<td><p>Number of threads for the kernel launch in X dimension.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__block_dim_y</span></code></p></td>
<td><p>Number of threads for the kernel launch in Y dimension.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__block_dim_z</span></code></p></td>
<td><p>Number of threads for the kernel launch in Z dimension.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__block_size</span></code></p></td>
<td><p>Total number of threads per block for the kernel launch.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__cluster_dim_x</span></code></p></td>
<td><p>Number of clusters for the kernel launch in X dimension.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__cluster_dim_y</span></code></p></td>
<td><p>Number of clusters for the kernel launch in Y dimension.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__cluster_dim_z</span></code></p></td>
<td><p>Number of clusters for the kernel launch in Z dimension.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__cluster_max_active</span></code></p></td>
<td><p>Maximum number of clusters that can co-exist on the target device. The runtime environment may affect how the hardware schedules the clusters, so the calculated occupancy is not guaranteed to be achievable.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__cluster_max_potential_size</span></code></p></td>
<td><p>Largest valid cluster size for the kernel function and launch configuration.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__cluster_scheduling_policy</span></code></p></td>
<td><p>Cluster scheduling policy.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__context_id</span></code></p></td>
<td><p>CUDA context id for the kernel launch (id of the primary context if launch was on a green context).</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__device_id</span></code></p></td>
<td><p>CUDA device id for the kernel launch.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__func_cache_config</span></code></p></td>
<td><p>On devices where the L1 cache and shared memory use the same hardware resources, this is the preferred cache configuration for the CUDA function. The runtime will use the requested configuration if possible, but it is free to choose a different configuration if required.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__function_pcs</span></code></p></td>
<td><p>Kernel function entry PCs.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__graph_contains_device_launch</span></code></p></td>
<td><p>Set to 1 if any node in the profiled graph can launch a CUDA device graph.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__graph_is_device_launchable</span></code></p></td>
<td><p>Set to 1 if the profiled graph was device-launchable.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__green_context_id</span></code></p></td>
<td><p>CUDA context id of the green context for the kernel launch (if applicable).</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__grid_dim_x</span></code></p></td>
<td><p>Number of blocks for the kernel launch in X dimension.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__grid_dim_y</span></code></p></td>
<td><p>Number of blocks for the kernel launch in Y dimension.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__grid_dim_z</span></code></p></td>
<td><p>Number of blocks for the kernel launch in Z dimension.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__grid_size</span></code></p></td>
<td><p>Total number of blocks for the kernel launch.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_cluster_gpu_pct</span></code></p></td>
<td><p>Overall GPU occupancy due to clusters.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_cluster_pct</span></code></p></td>
<td><p>The ratio of active blocks to the max possible active blocks due to clusters.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_limit_blocks</span></code></p></td>
<td><p>Occupancy limit due to maximum number of blocks managable per SM.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_limit_registers</span></code></p></td>
<td><p>Occupancy limit due to register usage.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_limit_shared_mem</span></code></p></td>
<td><p>Occupancy limit due to shared memory usage.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_limit_warps</span></code></p></td>
<td><p>Occupancy limit due to block size.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_per_block_size</span></code></p></td>
<td><p>Number of active warps for given block size.</p>
<p>Instance values map from number of warps (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_per_cluster_size</span></code></p></td>
<td><p>Number of active clusters for given cluster size.</p>
<p>Instance values map from number of clusters (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_per_register_count</span></code></p></td>
<td><p>Number of active warps for given register count.</p>
<p>Instance values map from number of warps (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__occupancy_per_shared_mem_size</span></code></p></td>
<td><p>Number of active warps for given shared memory size.</p>
<p>Instance values map from number of warps (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__registers_per_thread</span></code></p></td>
<td><p>Number of registers allocated per thread.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__registers_per_thread_allocated</span></code></p></td>
<td><p>Number of registers allocated per thread.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__shared_mem_config_size</span></code></p></td>
<td><p>Shared memory size configured for the kernel launch. The size depends on the static, dynamic, and driver shared memory requirements as well as the specified or platform-determined configuration size.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__shared_mem_per_block_allocated</span></code></p></td>
<td><p>Allocated shared memory size per block.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__shared_mem_per_block_driver</span></code></p></td>
<td><p>Shared memory size per block, allocated for the CUDA driver.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__shared_mem_per_block_dynamic</span></code></p></td>
<td><p>Dynamic shared memory size per block, allocated for the kernel.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__shared_mem_per_block_static</span></code></p></td>
<td><p>Static shared memory size per block, allocated for the kernel.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__sm_count</span></code></p></td>
<td><p>Number of SMs utilized in the launch.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__stream_id</span></code></p></td>
<td><p>CUDA stream id for the kernel launch.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__sub_launch_name</span></code></p></td>
<td><p>Name of each sub-launch for range-like results.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__thread_count</span></code></p></td>
<td><p>Total number of threads across all blocks for the kernel launch.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__uses_cdp</span></code></p></td>
<td><p>Set to 1 if any function object in the launched workload can use CUDA dynamic parallelism.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">launch__uses_green_context</span></code></p></td>
<td><p>Set to 1 if launch was on a green context.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">launch__waves_per_multiprocessor</span></code></p></td>
<td><p>Number of waves per SM. Partial waves can lead to tail effects where some SMs become idle while others still have pending work to complete.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-nvlink-topology">NVLink Topology Metrics</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id28">
<caption><span class="caption-text">NVLink Topology Metrics</span><a class="headerlink" href="#id28" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__bandwidth</span></code></p></td>
<td><p>Link bandwidth in bytes/s.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (double).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__count_logical</span></code></p></td>
<td><p>Total number of logical NVLinks.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__count_physical</span></code></p></td>
<td><p>Total number of physical links.</p>
<p>Instance values map from physical nvlink device ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__destination_ports</span></code></p></td>
<td><p>Destination port numbers (as strings).</p>
<p>Instance values map from logical nvlink ID (uint64) to comma-separated list of port numbers (string).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__dev0Id</span></code></p></td>
<td><p>ID of the first connected device.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__dev0type</span></code></p></td>
<td><p>Type of the first connected device.</p>
<p>Instance values map from logical nvlink ID (uint64) to values [1=GPU, 2=CPU] (uint64).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__dev1Id</span></code></p></td>
<td><p>ID of the second connected device.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__dev1type</span></code></p></td>
<td><p>Type of the second connected device.</p>
<p>Instance values map from logical nvlink ID (uint64) to values [1=GPU, 2=CPU] (uint64).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__dev_display_name_all</span></code></p></td>
<td><p>Device display name.</p>
<p>Instance values map from logical nvlink device ID (uint64) to value (string).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__enabled_mask</span></code></p></td>
<td><p>NVLink enablement mask, per device.</p>
<p>Instance values map from physical nvlink device ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__is_direct_link</span></code></p></td>
<td><p>Indicates, per NVLink, if the link is direct.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__is_nvswitch_connected</span></code></p></td>
<td><p>Indicates if NVSwitch is connected.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__max_count</span></code></p></td>
<td><p>Maximum number of NVLinks.</p>
<p>Instance values map from physical nvlink device ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__peer_access</span></code></p></td>
<td><p>Indicates if peer access is supported.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__peer_atomic</span></code></p></td>
<td><p>Indicates if peer atomics are supported.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__source_ports</span></code></p></td>
<td><p>Source port numbers (as strings).</p>
<p>Instance values map from logical nvlink ID (uint64) to comma-separated list of port numbers (string).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__system_access</span></code></p></td>
<td><p>Indicates if system access is supported.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (uint64).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">nvlink__system_atomic</span></code></p></td>
<td><p>Indicates if system atomics are supported.</p>
<p>Instance values map from logical nvlink ID (uint64) to value (uint64).</p>
</td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-numa-topology">NUMA Topology Metrics</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id29">
<caption><span class="caption-text">NUMA Topology Metrics</span><a class="headerlink" href="#id29" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">numa__cpu_affinity</span></code></p></td>
<td><p>CPU affinity for each device.</p>
<p>Instance values map from device ID (uint64) to comma-separated values (string).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">numa__dev_display_name_all</span></code></p></td>
<td><p>Device display names for all devices.</p>
<p>Instance values map from device ID (uint64) to comma-separated values (string).</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">numa__id_cpu</span></code></p></td>
<td><p>NUMA ID of the nearest CPU for each device.</p>
<p>Instance values map from device ID (uint64) to comma-separated values (string).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">numa__id_memory</span></code></p></td>
<td><p>NUMA ID of the nearest memory for each device.</p>
<p>Instance values map from device ID (uint64) to comma-separated values (string).</p>
</td>
</tr>
</tbody>
</table>
</div></blockquote>
<p class="rubric-h3 rubric" id="device-attributes">Device Attributes</p>
<p><code class="docutils literal notranslate"><span class="pre">device__attribute_*</span></code> metrics represent <a class="reference external" href="https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g49e2f8c2c0bd6fe264f2fc970912e5cd">CUDA device attributes</a>. Collecting them does not require an additional kernel replay pass, as their value is available from the CUDA driver for each CUDA device.</p>
<p>See below for custom <code class="docutils literal notranslate"><span class="pre">device__attribute_*</span></code> metrics.</p>
<blockquote>
<div><table class="colwidths-auto table-no-stripes docutils align-default">
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_architecture</span></code></p></td>
<td><p>Chip architecture of the CUDA device.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_confidential_computing_mode</span></code></p></td>
<td><p>Confidential computing mode.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_device_index</span></code></p></td>
<td><p>Device index.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_display_name</span></code></p></td>
<td><p>Product name of the CUDA device.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_fb_bus_width</span></code></p></td>
<td><p>Frame buffer bus width.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_fbp_count</span></code></p></td>
<td><p>Total number of frame buffer partitions.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_implementation</span></code></p></td>
<td><p>Chip implementation of the CUDA device.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_l2s_count</span></code></p></td>
<td><p>Total number of Level 2 cache slices.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_limits_max_cta_per_sm</span></code></p></td>
<td><p>Maximum number of CTA per SM.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_max_gpu_frequency_khz</span></code></p></td>
<td><p>Maximum GPU frequency in kilohertz.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_max_ipc_per_multiprocessor</span></code></p></td>
<td><p>Maximum number of instructions per clock per multiprocessor.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_max_ipc_per_scheduler</span></code></p></td>
<td><p>Maximum number of instructions per clock per scheduler.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_max_mem_frequency_khz</span></code></p></td>
<td><p>Peak memory frequency in kilohertz.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_max_registers_per_thread</span></code></p></td>
<td><p>Maximum number of registers available per thread.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_max_warps_per_multiprocessor</span></code></p></td>
<td><p>Maximum number of warps per multiprocessor.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_max_warps_per_scheduler</span></code></p></td>
<td><p>Maximum number of warps per scheduler.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_num_l2s_per_fbp</span></code></p></td>
<td><p>Number of Level 2 cache slices per frame buffer partition.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_num_schedulers_per_multiprocessor</span></code></p></td>
<td><p>Number of schedulers per multiprocessor.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_num_tex_per_multiprocessor</span></code></p></td>
<td><p>Number of TEX unit per multiprocessor.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">device__attribute_sass_level</span></code></p></td>
<td><p>SASS level.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-warp-stall-reasons">Warp Stall Reasons</p>
</div></blockquote>
<p>Collected using warp scheduler state sampling. They are incremented regardless if the scheduler issued an instruction in the same cycle or not. These metrics have instance values mapping from the function address (uint64) to the number of samples (uint64).</p>
<blockquote>
<div><table class="colwidths-auto table-no-stripes docutils align-default" id="id30">
<caption><span class="caption-text">Warp Stall Reasons</span><a class="headerlink" href="#id30" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_barrier</span></code></p></td>
<td><p>Warp was stalled waiting for sibling warps at a CTA barrier. A high number of warps waiting at a barrier is commonly caused by diverging code paths before a barrier. This causes some warps to wait a long time until other warps reach the synchronization point. Whenever possible, try to divide up the work into blocks of uniform workloads. If the block size is 512 threads or greater, consider splitting it into smaller groups. This can increase eligible warps without affecting occupancy, unless shared memory becomes a new occupancy limiter. Also, try to identify which barrier instruction causes the most stalls, and optimize the code executed before that synchronization point first.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_branch_resolving</span></code></p></td>
<td><p>Warp was stalled waiting for a branch target to be computed, and the warp program counter to be updated. To reduce the number of stalled cycles, consider using fewer jump/branch operations and reduce control flow divergence, e.g. by reducing or coalescing conditionals in your code. See also the related No Instructions state.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_dispatch_stall</span></code></p></td>
<td><p>Warp was stalled waiting on a dispatch stall. A warp stalled during dispatch has an instruction ready to issue, but the dispatcher holds back issuing the warp due to other conflicts or events.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_drain</span></code></p></td>
<td><p>Warp was stalled after EXIT waiting for all outstanding memory operations to complete so that warp’s resources can be freed. A high number of stalls due to draining warps typically occurs when a lot of data is written to memory towards the end of a kernel. Make sure the memory access patterns of these store operations are optimal for the target architecture and consider parallelized data reduction, if applicable.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_imc_miss</span></code></p></td>
<td><p>Warp was stalled waiting for an immediate constant cache (IMC) miss. A read from constant memory costs one memory read from device memory only on a cache miss; otherwise, it just costs one read from the constant cache. Immediate constants are encoded into the SASS instruction as ‘c[bank][offset]’. Accesses to different addresses by threads within a warp are serialized, thus the cost scales linearly with the number of unique addresses read by all threads within a warp. As such, the constant cache is best when threads in the same warp access only a few distinct locations. If all threads of a warp access the same location, then constant memory can be as fast as a register access.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_lg_throttle</span></code></p></td>
<td><p>Warp was stalled waiting for the L1 instruction queue for local and global (LG) memory operations to be not full. Typically, this stall occurs only when executing local or global memory instructions extremely frequently. Avoid redundant global memory accesses. Try to avoid using thread-local memory by checking if dynamically indexed arrays are declared in local scope, of if the kernel has excessive register pressure causing by spills. If applicable, consider combining multiple lower-width memory operations into fewer wider memory operations and try interleaving memory operations and math instructions.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_long_scoreboard</span></code></p></td>
<td><p>Warp was stalled waiting for a scoreboard dependency on a L1TEX (local, global, surface, texture) operation. Find the instruction producing the data being waited upon to identify the culprit. To reduce the number of cycles waiting on L1TEX data accesses verify the memory access patterns are optimal for the target architecture, attempt to increase cache hit rates by increasing data locality (coalescing), or by changing the cache configuration. Consider moving frequently used data to shared memory.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_math_pipe_throttle</span></code></p></td>
<td><p>Warp was stalled waiting for the execution pipe to be available. This stall occurs when all active warps execute their next instruction on a specific, oversubscribed math pipeline. Try to increase the number of active warps to hide the existent latency or try changing the instruction mix to utilize all available pipelines in a more balanced way.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_membar</span></code></p></td>
<td><p>Warp was stalled waiting on a memory barrier. Avoid executing any unnecessary memory barriers and assure that any outstanding memory operations are fully optimized for the target architecture.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_mio_throttle</span></code></p></td>
<td><p>Warp was stalled waiting for the MIO (memory input/output) instruction queue to be not full. This stall reason is high in cases of extreme utilization of the MIO pipelines, which include special math instructions, dynamic branches, as well as shared memory instructions. When caused by shared memory accesses, trying to use fewer but wider loads can reduce pipeline pressure.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_misc</span></code></p></td>
<td><p>Warp was stalled for a miscellaneous hardware reason.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_no_instructions</span></code></p></td>
<td><p>Warp was stalled waiting to be selected to fetch an instruction or waiting on an instruction cache miss. A high number of warps not having an instruction fetched is typical for very short kernels with less than one full wave of work in the grid. Excessively jumping across large blocks of assembly code can also lead to more warps stalled for this reason, if this causes misses in the instruction cache. See also the related Branch Resolving state.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_not_selected</span></code></p></td>
<td><p>Warp was stalled waiting for the micro scheduler to select the warp to issue. Not selected warps are eligible warps that were not picked by the scheduler to issue that cycle as another warp was selected. A high number of not selected warps typically means you have sufficient warps to cover warp latencies and you may consider reducing the number of active warps to possibly increase cache coherence and data locality.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_selected</span></code></p></td>
<td><p>Warp was selected by the micro scheduler and issued an instruction.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_short_scoreboard</span></code></p></td>
<td><p>Warp was stalled waiting for a scoreboard dependency on a MIO (memory input/output) operation (not to L1TEX). The primary reason for a high number of stalls due to short scoreboards is typically memory operations to shared memory. Other reasons include frequent execution of special math instructions (e.g. MUFU) or dynamic branching (e.g. BRX, JMX). Consult the Memory Workload Analysis section to verify if there are shared memory operations and reduce bank conflicts, if reported. Assigning frequently accessed values to variables can assist the compiler in using low-latency registers instead of direct memory accesses.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_sleeping</span></code></p></td>
<td><p>Warp was stalled due to all threads in the warp being in the blocked, yielded, or sleep state. Reduce the number of executed NANOSLEEP instructions, lower the specified time delay, and attempt to group threads in a way that multiple threads in a warp sleep at the same time.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_tex_throttle</span></code></p></td>
<td><p>Warp was stalled waiting for the L1 instruction queue for texture operations to be not full. This stall reason is high in cases of extreme utilization of the L1TEX pipeline. Try issuing fewer texture fetches, surface loads, surface stores, or decoupled math operations. If applicable, consider combining multiple lower-width memory operations into fewer wider memory operations and try interleaving memory operations and math instructions. Consider converting texture lookups or surface loads into global memory lookups. Texture can accept four threads’ requests per cycle, whereas global accepts 32 threads.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_wait</span></code></p></td>
<td><p>Warp was stalled waiting on a fixed latency execution dependency. Typically, this stall reason should be very low and only shows up as a top contributor in already highly optimized kernels. Try to hide the corresponding instruction latencies by increasing the number of active warps, restructuring the code or unrolling loops. Furthermore, consider switching to lower-latency instructions, e.g. by making use of fast math compiler options.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-warp-stall-reasons-not-issued">Warp Stall Reasons (Not Issued)</p>
</div></blockquote>
<p>Collected using warp scheduler state sampling. They are incremented only on cycles in which the warp scheduler issued no instruction. These metrics have instance values mapping from the function address (uint64) to the number of samples (uint64).</p>
<blockquote>
<div><table class="colwidths-auto table-no-stripes docutils align-default" id="id31">
<caption><span class="caption-text">Warp Stall Reasons (Not Issued)</span><a class="headerlink" href="#id31" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_barrier_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for sibling warps at a CTA barrier. A high number of warps waiting at a barrier is commonly caused by diverging code paths before a barrier. This causes some warps to wait a long time until other warps reach the synchronization point. Whenever possible, try to divide up the work into blocks of uniform workloads. If the block size is 512 threads or greater, consider splitting it into smaller groups. This can increase eligible warps without affecting occupancy, unless shared memory becomes a new occupancy limiter. Also, try to identify which barrier instruction causes the most stalls, and optimize the code executed before that synchronization point first.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_branch_resolving_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for a branch target to be computed, and the warp program counter to be updated. To reduce the number of stalled cycles, consider using fewer jump/branch operations and reduce control flow divergence, e.g. by reducing or coalescing conditionals in your code. See also the related No Instructions state.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_dispatch_stall_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting on a dispatch stall. A warp stalled during dispatch has an instruction ready to issue, but the dispatcher holds back issuing the warp due to other conflicts or events.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_drain_not_issued</span></code></p></td>
<td><p>Warp was stalled after EXIT waiting for all memory operations to complete so that warp resources can be freed. A high number of stalls due to draining warps typically occurs when a lot of data is written to memory towards the end of a kernel. Make sure the memory access patterns of these store operations are optimal for the target architecture and consider parallelized data reduction, if applicable.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_imc_miss_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for an immediate constant cache (IMC) miss. A read from constant memory costs one memory read from device memory only on a cache miss; otherwise, it just costs one read from the constant cache. Accesses to different addresses by threads within a warp are serialized, thus the cost scales linearly with the number of unique addresses read by all threads within a warp. As such, the constant cache is best when threads in the same warp access only a few distinct locations. If all threads of a warp access the same location, then constant memory can be as fast as a register access.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_lg_throttle_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for the L1 instruction queue for local and global (LG) memory operations to be not full. Typically, this stall occurs only when executing local or global memory instructions extremely frequently. Avoid redundant global memory accesses. Try to avoid using thread-local memory by checking if dynamically indexed arrays are declared in local scope, of if the kernel has excessive register pressure causing by spills. If applicable, consider combining multiple lower-width memory operations into fewer wider memory operations and try interleaving memory operations and math instructions.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_long_scoreboard_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for a scoreboard dependency on a L1TEX (local, global, surface, texture) operation. Find the instruction producing the data being waited upon to identify the culprit. To reduce the number of cycles waiting on L1TEX data accesses verify the memory access patterns are optimal for the target architecture, attempt to increase cache hit rates by increasing data locality (coalescing), or by changing the cache configuration. Consider moving frequently used data to shared memory.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_math_pipe_throttle_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for the execution pipe to be available. This stall occurs when all active warps execute their next instruction on a specific, oversubscribed math pipeline. Try to increase the number of active warps to hide the existent latency or try changing the instruction mix to utilize all available pipelines in a more balanced way.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_membar_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting on a memory barrier. Avoid executing any unnecessary memory barriers and assure that any outstanding memory operations are fully optimized for the target architecture.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_mio_throttle_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for the MIO (memory input/output) instruction queue to be not full. This stall reason is high in cases of extreme utilization of the MIO pipelines, which include special math instructions, dynamic branches, as well as shared memory instructions. When caused by shared memory accesses, trying to use fewer but wider loads can reduce pipeline pressure.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_misc_not_issued</span></code></p></td>
<td><p>Warp was stalled for a miscellaneous hardware reason.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_no_instructions_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting to be selected to fetch an instruction or waiting on an instruction cache miss. A high number of warps not having an instruction fetched is typical for very short kernels with less than one full wave of work in the grid. Excessively jumping across large blocks of assembly code can also lead to more warps stalled for this reason, if this causes misses in the instruction cache. See also the related Branch Resolving state.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_not_selected_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for the micro scheduler to select the warp to issue. Not selected warps are eligible warps that were not picked by the scheduler to issue that cycle as another warp was selected. A high number of not selected warps typically means you have sufficient warps to cover warp latencies and you may consider reducing the number of active warps to possibly increase cache coherence and data locality.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_selected_not_issued</span></code></p></td>
<td><p>Warp was selected by the micro scheduler and issued an instruction.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_short_scoreboard_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for a scoreboard dependency on a MIO (memory input/output) operation (not to L1TEX). The primary reason for a high number of stalls due to short scoreboards is typically memory operations to shared memory. Other reasons include frequent execution of special math instructions (e.g. MUFU) or dynamic branching (e.g. BRX, JMX). Consult the Memory Workload Analysis section to verify if there are shared memory operations and reduce bank conflicts, if reported. Assigning frequently accessed values to variables can assist the compiler in using low-latency registers instead of direct memory accesses.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_sleeping_not_issued</span></code></p></td>
<td><p>Warp was stalled due to all threads in the warp being in the blocked, yielded, or sleep state. Reduce the number of executed NANOSLEEP instructions, lower the specified time delay, and attempt to group threads in a way that multiple threads in a warp sleep at the same time.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_tex_throttle_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting for the L1 instruction queue for texture operations to be not full. This stall reason is high in cases of extreme utilization of the L1TEX pipeline. Try issuing fewer texture fetches, surface loads, surface stores, or decoupled math operations. If applicable, consider combining multiple lower-width memory operations into fewer wider memory operations and try interleaving memory operations and math instructions. Consider converting texture lookups or surface loads into global memory lookups. Texture can accept four threads’ requests per cycle, whereas global accepts 32 threads.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_warps_issue_stalled_wait_not_issued</span></code></p></td>
<td><p>Warp was stalled waiting on a fixed latency execution dependency. Typically, this stall reason should be very low and only shows up as a top contributor in already highly optimized kernels. Try to hide the corresponding instruction latencies by increasing the number of active warps, restructuring the code or unrolling loops. Furthermore, consider switching to lower-latency instructions, e.g. by making use of fast math compiler options.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-source">Source Metrics</p>
</div></blockquote>
<p>Most are collected using SASS-patching. These metrics have instance values mapping from function address (uint64) to associated values (uint64). Metrics <code class="docutils literal notranslate"><span class="pre">memory_[access_]type</span></code> map to string values.</p>
<blockquote>
<div><table class="colwidths-auto table-no-stripes docutils align-default" id="id32">
<caption><span class="caption-text">Source Metrics</span><a class="headerlink" href="#id32" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">branch_inst_executed</span></code></p></td>
<td><p>Number of unique branch targets assigned to the instruction, including both divergent and uniform branches.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">derived__avg_thread_executed</span></code></p></td>
<td><p>Average number of thread-level executed instructions per warp (regardless of their predicate). Computed as: thread_inst_executed / inst_executed</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">derived__avg_thread_executed_true</span></code></p></td>
<td><p>Average number of predicated-on thread-level executed instructions per warp. Computed as: thread_inst_executed_true / inst_executed</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">derived__memory_l1_conflicts_shared_nway</span></code></p></td>
<td><p>Average N-way conflict in L1 per shared memory instruction. A 1-way access has no conflicts and resolves in a single pass. Computed as: memory_l1_wavefronts_shared / inst_executed</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">derived__memory_l1_wavefronts_shared_excessive</span></code></p></td>
<td><p>Excessive number of wavefronts in L1 from shared memory instructions, because not all not predicated-off threads performed the operation.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">derived__memory_l2_theoretical_sectors_global_excessive</span></code></p></td>
<td><p>Excessive theoretical number of sectors requested in L2 from global memory instructions, because not all not predicated-off threads performed the operation.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">inst_executed</span></code></p></td>
<td><p>Number of warp-level executed instructions, ignoring instruction predicates. Warp-level means the values increased by one per individual warp executing the instruction, independent of the number of participating threads within each warp.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">memory_access_size_type</span></code></p></td>
<td><p>The size of the memory access, in bits.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">memory_access_type</span></code></p></td>
<td><p>The type of memory access (e.g. load or store).</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">memory_l1_tag_requests_global</span></code></p></td>
<td><p>Number of L1 tag requests generated by global memory instructions.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">memory_l1_wavefronts_shared</span></code></p></td>
<td><p>Number of wavefronts in L1 from shared memory instructions.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">memory_l1_wavefronts_shared_ideal</span></code></p></td>
<td><p>Ideal number of wavefronts in L1 from shared memory instructions, assuming each not predicated-off thread performed the operation.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">memory_l2_theoretical_sectors_global</span></code></p></td>
<td><p>Theoretical number of sectors requested in L2 from global memory instructions.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">memory_l2_theoretical_sectors_global_ideal</span></code></p></td>
<td><p>Ideal number of sectors requested in L2 from global memory instructions, assuming each not predicated-off thread performed the operation.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">memory_l2_theoretical_sectors_local</span></code></p></td>
<td><p>Theoretical number of sectors requested in L2 from local memory instructions.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">memory_type</span></code></p></td>
<td><p>The accessed address space (global/local/shared).</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__branch_targets_threads_divergent</span></code></p></td>
<td><p>Number of divergent branch targets, including fallthrough. Incremented only when there are two or more active threads with divergent targets.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__branch_targets_threads_uniform</span></code></p></td>
<td><p>Number of uniform branch execution, including fallthrough, where all active threads selected the same branch target.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_sample_count</span></code></p></td>
<td><p>Number of collected warp state samples per program counter. This metric is collected using warp sampling.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">thread_inst_executed</span></code></p></td>
<td><p>Number of thread-level executed instructions, regardless of predicate presence or evaluation.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">thread_inst_executed_true</span></code></p></td>
<td><p>Number of thread-level executed instructions, where the instruction predicate evaluated to true, or no predicate was given.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-memdesc">L2 Cache Eviction Metrics</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id33">
<caption><span class="caption-text">L2 Cache Eviction Metrics</span><a class="headerlink" href="#id33" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__sass_inst_executed_memdesc_explicit_evict_type</span></code></p></td>
<td><p>L2 cache eviction policy types.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_first</span></code></p></td>
<td><p>Number of warp-level executed instructions with L2 cache eviction hit property ‘first’.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_last</span></code></p></td>
<td><p>Number of warp-level executed instructions with L2 cache eviction hit property ‘last’.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_normal</span></code></p></td>
<td><p>Number of warp-level executed instructions with L2 cache eviction hit property ‘normal’.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_normal_demote</span></code></p></td>
<td><p>Number of warp-level executed instructions with L2 cache eviction hit property ‘normal demote’.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__sass_inst_executed_memdesc_explicit_missprop_evict_first</span></code></p></td>
<td><p>Number of warp-level executed instructions with L2 cache eviction miss property ‘first’.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__sass_inst_executed_memdesc_explicit_missprop_evict_normal</span></code></p></td>
<td><p>Number of warp-level executed instructions with L2 cache eviction miss property ‘normal’.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-sass-per-opcode">Instructions Per Opcode Metrics</p>
</div></blockquote>
<p>Collected using SASS-patching. These metrics have instance values mapping from the SASS opcode (string) to the number of executions (uint64).</p>
<blockquote>
<div><table class="colwidths-auto table-no-stripes docutils align-default" id="id34">
<caption><span class="caption-text">Instructions Per Opcode Metrics</span><a class="headerlink" href="#id34" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">sass__inst_executed_per_opcode</span></code></p></td>
<td><p>Number of warp-level executed instructions, instanced by basic SASS opcode.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">sass__inst_executed_per_opcode_with_modifier_all</span></code></p></td>
<td><p>Number of warp-level executed instructions, instanced by all SASS opcode modifiers.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">sass__inst_executed_per_opcode_with_modifier_selective</span></code></p></td>
<td><p>Number of warp-level executed instructions, instanced by selective SASS opcode modifiers.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">sass__thread_inst_executed_true_per_opcode</span></code></p></td>
<td><p>Number of thread-level executed instructions, instanced by basic SASS opcode.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">sass__thread_inst_executed_true_per_opcode_with_modifier_all</span></code></p></td>
<td><p>Number of thread-level executed instructions, instanced by all SASS opcode modifiers.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">sass__thread_inst_executed_true_per_opcode_with_modifier_selective</span></code></p></td>
<td><p>Number of thread-level executed instructions, instanced by selective SASS opcode modifiers.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-groups">Metric Groups</p>
<table class="colwidths-auto table-no-stripes docutils align-default" id="id35">
<caption><span class="caption-text">Metric Groups</span><a class="headerlink" href="#id35" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">group:memory__chart</span></code></p></td>
<td><p>Group of metrics for the workload analysis chart.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">group:memory__dram_table</span></code></p></td>
<td><p>Group of metrics for the device memory workload analysis table.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">group:memory__first_level_cache_table</span></code></p></td>
<td><p>Group of metrics for the L1/TEX cache workload analysis table.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">group:memory__l2_cache_evict_policy_table</span></code></p></td>
<td><p>Group of metrics for the L2 cache eviction policies table.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">group:memory__l2_cache_table</span></code></p></td>
<td><p>Group of metrics for the L2 cache workload analysis table.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">group:memory__shared_table</span></code></p></td>
<td><p>Group of metrics for the shared memory workload analysis table.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">group:smsp__pcsamp_warp_stall_reasons</span></code></p></td>
<td><p>Group of metrics for the number of samples from the warp sampler per program location.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">group:smsp__pcsamp_warp_stall_reasons_not_issued</span></code></p></td>
<td><p>Group of metrics for the number of samples from the warp sampler per program location on cycles the warp scheduler issued no instructions.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-ref-profiler">Profiler Metrics</p>
</div></blockquote>
<p>Metrics generated by the tool itself to inform about statistics or problems during profiling.</p>
<blockquote>
<div><table class="colwidths-auto table-no-stripes docutils align-default" id="id36">
<caption><span class="caption-text">Profiler Metrics</span><a class="headerlink" href="#id36" title="Permalink to this table"></a></caption>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__perfworks_session_reuse</span></code></p></td>
<td><p>Indicates if the PerfWorks session was reused between results.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__pmsampler_buffer_size_bytes</span></code></p></td>
<td><p>Buffer size in bytes per pass group used for PM sampling.</p>
<p>Instance values map from pass group to bytes.</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__pmsampler_ctxsw_*</span></code></p></td>
<td><p>GPU context switch states over time during PM sampling for a specific pass group.</p>
<p>Instance values map from timestamp to context state (1 - enabled, 0 - disabled).</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__pmsampler_dropped_samples</span></code></p></td>
<td><p>Number of samples dropped per pass group during PM sampling due to insufficient buffer size.</p>
<p>Instance values map from pass group to samples.</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__pmsampler_interval_cycles</span></code></p></td>
<td><p>Sampling interval in cycles per pass group used for PM sampling, or zero if time-based interval was used.</p>
<p>Instance values map from pass group to cycles.</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__pmsampler_interval_time</span></code></p></td>
<td><p>Sampling interval in nanoseconds per pass group used for PM sampling, or zero if cycle-based interval was used.</p>
<p>Instance values map from pass group to nanoseconds.</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__pmsampler_merged_samples</span></code></p></td>
<td><p>Number of samples merged per pass group during PM sampling due to HW back pressure while streaming results.</p>
<p>Instance values map from pass group to samples.</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__pmsampler_pass_groups</span></code></p></td>
<td><p>Number of pass groups used for PM sampling.</p>
<p>Instance values map from pass group to comma-separated list of metrics collected in this pass.</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__replayer_passes</span></code></p></td>
<td><p>Number of passes the result was replayed for profiling across all experiments.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">profiler__replayer_passes_type_warmup</span></code></p></td>
<td><p>Number of passes the result was replayed to warmup the GPU for profiling.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_aggregated_passes</span></code></p></td>
<td><p>Number of passes required for statistical warp stall sampling.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_buffer_size_bytes</span></code></p></td>
<td><p>Buffer size in bytes for statistical warp stall sampling.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_dropped_bytes</span></code></p></td>
<td><p>Bytes dropped during statistical warp stall sampling due to insufficient buffer size.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_interval</span></code></p></td>
<td><p>Interval number for warp stall sampling.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">smsp__pcsamp_interval_cycles</span></code></p></td>
<td><p>Interval cycles for statistical warp stall sampling.</p></td>
</tr>
</tbody>
</table>
</div></blockquote>
</section>
<section id="sampling">
<h2><span class="section-number">2.5. </span>Sampling<a class="headerlink" href="#sampling" title="Permalink to this headline"></a></h2>
<p>NVIDIA Nsight Compute can collect certain performance data via sampling at fixed intervals.</p>
<section id="pm-sampling">
<h3><span class="section-number">2.5.1. </span>PM Sampling<a class="headerlink" href="#pm-sampling" title="Permalink to this headline"></a></h3>
<p>NVIDIA Nsight Compute supports collecting many metrics by sampling the GPU’s performance monitors (PM) periodically at fixed intervals. The resulting metrics are <a class="reference external" href="index.html#instanced-metrics">instanced</a>, with each sample being composed of its value and the (GPU) timestamp when it was collected. This allows the tool to visualize the data on a <a class="reference external" href="../NsightCompute/index.html#profiler-report-details-page">timeline</a> that helps you understand how the behavior of the profiled workload changes during its runtime.</p>
<p>Metrics collected with PM sampling have instance values mapping from their sample timestamp (in ns) to their sample value. When logically possible, the non-instanced value of the metric represents the aggregate across all instances. The aggregation operation (e.g. sum, average) depends on the metric structure.</p>
<p>A metric is collected using PM sampling in the following cases:</p>
<ul class="simple">
<li><p>The metric name has the <code class="docutils literal notranslate"><span class="pre">pmsampling:</span></code> <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">prefix</a>.</p></li>
<li><p>The metric name includes a valid <code class="docutils literal notranslate"><span class="pre">Triage</span></code> group.</p></li>
<li><p>The metric is requested in a section’s <code class="docutils literal notranslate"><span class="pre">Timeline</span></code> field. Prefixing the metric with <code class="docutils literal notranslate"><span class="pre">pmsampling:</span></code> is still recommended in this case to avoid conflicts with profiler metrics of the same name collected e.g. by other sections.</p></li>
</ul>
<table class="docutils align-default" id="id37">
<caption><span class="caption-text">Supported Architectures for PM Sampling</span><a class="headerlink" href="#id37" title="Permalink to this table"></a></caption>
<colgroup>
<col style="width: 33%" />
<col style="width: 33%" />
<col style="width: 33%" />
</colgroup>
<thead>
<tr class="row-odd"><th class="head"><p>Architecture</p></th>
<th class="head"><p>Support</p></th>
<th class="head"><p>Sampling Intervals</p></th>
</tr>
</thead>
<tbody>
<tr class="row-even"><td><p>Volta and earlier</p></td>
<td><p>Not supported</p></td>
<td><p>n/a</p></td>
</tr>
<tr class="row-odd"><td><p>TU10x-GA100</p></td>
<td><p>Supported</p></td>
<td><p>>= 20000 cycles</p></td>
</tr>
<tr class="row-even"><td><p>GA10x and later</p></td>
<td><p>Supported</p></td>
<td><p>>= 1000 ns</p></td>
</tr>
</tbody>
</table>
<p>PM sampling is supported on all platforms except vGPU. See below for further limitations that apply to <a class="reference external" href="index.html#ctx-switch-trace">context switch trace</a>.
You can <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">query</a> the list of metrics available to PM sampling using the <code class="docutils literal notranslate"><span class="pre">--query-metrics-collection</span> <span class="pre">pmsampling</span></code> option.
Note though that while all listed metrics are available to the PM sampler, only those requiring a single pass can be collected.</p>
<p class="title sectiontitle rubric" id="ctx-switch-trace">Context Switch Trace</p>
<p>Since this data collection samples across the entire GPU device, the tool concurrently collects a <em>context switch trace</em>. The trace is stored as a separate, instanced <a class="reference external" href="index.html#profiler-metrics">metric</a>. It tracks when the context of interest was active and can be used to filter the sampling metric to only relevant instances and to better align metrics from multiple passes on the timeline. While it’s generally preferable to have this trace collected, it can be disabled using an <a class="reference external" href="../NsightComputeCli/index.html#environment-variables">environment variable</a>.</p>
<p>Note that context switch trace is not supported on Windows Subsystem for Linux (WSL), Multi-Instance GPU (MIG), within containers or on mobile platforms.</p>
<p class="title sectiontitle rubric" id="counter-domains">Counter Domains</p>
<p>PM sampling metrics are composed of one or more raw counter dependencies internally. If metrics in the same pass share such a dependency, it is only collected once.
Each counter is associated with a <em>counter domain</em>, which describes how and where in the hardware the counter is collected.
There is a limited number of counters in each domain that can be collected concurrently in the same pass, and the number may vary, depending on the selected counters.</p>
<p>Selecting counters from different domains has the possibility that more metric dependencies can be fit into the same pass.
Furthermore, some counters can be collected through different domains, and the domain may be chosen by the tool or the user.</p>
<p>When <a class="reference external" href="../NsightComputeCli/index.html#command-line-options-profile">querying</a> the PM sampling metric collection, the required and optional domains for a metric’s counter dependencies are shown.
E.g., for <code class="docutils literal notranslate"><span class="pre">l1tex__throughput</span> <span class="pre">gpu_sm_a,[gpu_sm_b,gpu_sm_c]</span></code>, the domain <code class="docutils literal notranslate"><span class="pre">gpu_sm_a</span></code> is required and one of the optional domains <code class="docutils literal notranslate"><span class="pre">[gpu_sm_b,gpu_sm_c]</span></code> must be chosen for this metric to be collectable.
Counter domains can only be selected explicitly in <a class="reference external" href="../CustomizationGuide/index.html#counter-domains">section files</a>, using one or more instances of the <code class="docutils literal notranslate"><span class="pre">CtrDomains:</span> <span class="pre">"<domain>"</span></code> field for PM sampling metrics.</p>
<p>Note that most users should be able to rely on the tool’s automatic selection of counter domains, or the pre-configured domains in section files.</p>
</section>
<section id="warp-sampling">
<span id="statistical-sampler"></span><h3><span class="section-number">2.5.2. </span>Warp Sampling<a class="headerlink" href="#warp-sampling" title="Permalink to this headline"></a></h3>
<p>NVIDIA Nsight Compute supports periodic sampling of the warp program counter and warp scheduler state. At a fixed interval of cycles, the sampler in each streaming multiprocessor selects an active warp and outputs the program counter and the warp scheduler state. The tool selects the minimum interval for the device. On small devices, this can be every 32 cycles. On larger chips with more multiprocessors, this may be 2048 cycles. The sampler selects a random active warp. On the same cycle the scheduler may select a different warp to issue.</p>
<p>The resulting metrics are correlated with the individual executed instructions but don’t have any time resolution.</p>
<p>See the <em>Warp Stall Reasons</em> tables in the <a class="reference external" href="index.html#metrics-reference">Metrics Reference</a> for a description of the individual warp scheduler states.</p>
</section>
</section>
<section id="reproducibility">
<h2><span class="section-number">2.6. </span>Reproducibility<a class="headerlink" href="#reproducibility" title="Permalink to this headline"></a></h2>
<p>In order to provide actionable and deterministic results across application runs, NVIDIA Nsight Compute applies various methods to adjust how metrics are collected. This includes <a class="reference external" href="index.html#serialization">serializing</a> kernel launches, <a class="reference external" href="index.html#cache-control">purging GPU caches</a> before each kernel replay or <a class="reference external" href="index.html#clock-control">adjusting GPU clocks</a>.</p>
<section id="serialization">
<h3><span class="section-number">2.6.1. </span>Serialization<a class="headerlink" href="#serialization" title="Permalink to this headline"></a></h3>
<p>NVIDIA Nsight Compute serializes kernel launches within the profiled application, potentially across multiple processes profiled by one or more instances of the tool at the same time.</p>
<p>Serialization across processes is necessary since for the collection of HW performance metrics, some GPU and driver objects can only be acquired by a single process at a time. To achieve this, the lock file <code class="docutils literal notranslate"><span class="pre">TMPDIR/nsight-compute-lock</span></code> is used. On Windows, <em>TMPDIR</em> is the path returned by the Windows <code class="docutils literal notranslate"><span class="pre">GetTempPath</span></code> API function. On other platforms, it is the path supplied by the first environment variable in the list <code class="docutils literal notranslate"><span class="pre">TMPDIR,</span> <span class="pre">TMP,</span> <span class="pre">TEMP,</span> <span class="pre">TEMPDIR</span></code>. If none of these is found, it’s <code class="docutils literal notranslate"><span class="pre">/var/nvidia</span></code> on QNX and <code class="docutils literal notranslate"><span class="pre">/tmp</span></code> otherwise.</p>
<p>Serialization within the process is required for most metrics to be mapped to the proper kernel. In addition, without serialization, performance metric values might vary widely if kernel execute concurrently on the same device.</p>
<p>It is currently not possible to disable this tool behavior. Refer to the <a class="reference external" href="index.html#faq">FAQ</a> entry on possible workarounds.</p>
</section>
<section id="clock-control">
<h3><span class="section-number">2.6.2. </span>Clock Control<a class="headerlink" href="#clock-control" title="Permalink to this headline"></a></h3>
<p>For many metrics, their value is directly influenced by the current GPU SM and memory clock frequencies. For example, if a kernel instance is profiled that has prior kernel executions in the application, the GPU might already be in a higher clocked state and the measured kernel duration, along with other metrics, will be affected. Likewise, if a kernel instance is the first kernel to be launched in the application, GPU clocks will regularly be lower. In addition, due to kernel replay, the metric value might depend on which replay pass it is collected in, as later passes would result in higher clock states.</p>
<p>To mitigate this non-determinism, NVIDIA Nsight Compute attempts to limit GPU clock frequencies to their <em>base</em> value. As a result, metric values are less impacted by the location of the kernel in the application, or by the number of the specific replay pass.</p>
<p>However, this behavior might be undesirable for analysis of the kernel, e.g. in cases where an external tool is used to fix clock frequencies, or where the behavior of the kernel within the application is analyzed. To solve this, users can adjust the <code class="docutils literal notranslate"><span class="pre">--clock-control</span></code> option to specify if any clock frequencies should be fixed by the tool.</p>
<p>Factors affecting Clock Control:</p>
<ul class="simple">
<li><p>Note that thermal throttling directed by the driver cannot be controlled by the tool and always overrides any selected options.</p></li>
<li><p>On mobile targets, e.g. L4T or QNX, there may be variations in profiling results due the inability for the tool to lock clocks. Using Nsight Compute’s <code class="docutils literal notranslate"><span class="pre">--clock-control</span></code> to set the GPU clocks will fail or will be silently ignored when profiling on a GPU partition.</p>
<ul>
<li><p>On L4T, you can use the jetson_clocks script to lock the clocks at their maximums during profiling.</p></li>
</ul>
</li>
<li><p>See the <a class="reference external" href="index.html#special-configurations">Special Configurations</a> section for MIG and vGPU clock control.</p></li>
</ul>
</section>
<section id="cache-control">
<h3><span class="section-number">2.6.3. </span>Cache Control<a class="headerlink" href="#cache-control" title="Permalink to this headline"></a></h3>
<p>As explained in <a class="reference external" href="index.html#kernel-replay">Kernel Replay</a>, the kernel might need to be replayed multiple times to collect all requested metrics. While NVIDIA Nsight Compute can save and restore the contents of GPU device memory accessed by the kernel for each pass, it cannot do the same for the contents of HW caches, such as e.g. the L1 and L2 cache.</p>
<p>This can have the effect that later replay passes might have better or worse performance than e.g. the first pass, as the caches could already be primed with the data last accessed by the kernel. Similarly, the values of HW performance counters collected by the first pass might depend on which kernels, if any, were executed prior to the measured kernel launch.</p>
<p>In order to make HW performance counter value more deterministic, NVIDIA Nsight Compute by default flushes all GPU caches before each replay pass. As a result, in each pass, the kernel will access a clean cache and the behavior will be as if the kernel was executed in complete isolation.</p>
<p>This behavior might be undesirable for performance analysis, especially if the measurement focuses on a kernel within a larger application execution, and if the collected data targets cache-centric metrics. In this case, you can use <code class="docutils literal notranslate"><span class="pre">--cache-control</span> <span class="pre">none</span></code> to disable flushing of any HW cache by the tool.</p>
</section>
<section id="persistence-mode">
<h3><span class="section-number">2.6.4. </span>Persistence Mode<a class="headerlink" href="#persistence-mode" title="Permalink to this headline"></a></h3>
<p>The NVIDIA kernel mode driver must be running and connected to a target GPU device before any user interactions with that device can take place. The driver behavior differs depending on the OS. Generally, on Linux, if the kernel mode driver is not already running or connected to a target GPU, the invocation of any program that attempts to interact with that GPU will transparently cause the driver to load and/or initialize the GPU. When all GPU clients terminate the driver will then deinitialize the GPU.</p>
<p>If <a class="reference external" href="https://docs.nvidia.com/deploy/driver-persistence/index.html">persistence mode</a> is not enabled (as part of the OS, or by the user), applications triggering GPU initialization may incur a short startup cost. In addition, on some configurations, there may also be a shutdown cost when the GPU is de-initialized at the end of the application.</p>
<p>It is recommended to enable persistence mode on applicable operating systems before profiling with NVIDIA Nsight Compute for more consistent application behavior.</p>
</section>
</section>
<section id="special-configurations">
<h2><span class="section-number">2.7. </span>Special Configurations<a class="headerlink" href="#special-configurations" title="Permalink to this headline"></a></h2>
<section id="multi-instance-gpu">
<span id="special-configurations-mig"></span><h3><span class="section-number">2.7.1. </span>Multi Instance GPU<a class="headerlink" href="#multi-instance-gpu" title="Permalink to this headline"></a></h3>
<p>Multi-Instance GPU (MIG) is a feature that allows a GPU to be partitioned into multiple CUDA devices. The partitioning is carried out on two levels: First, a GPU can be split into one or multiple GPU Instances. Each GPU Instance claims ownership of one or more streaming multiprocessors (SM), a subset of the overall GPU memory, and possibly other GPU resources, such as the video encoders/decoders. Second, each GPU Instance can be further partitioned into one or more Compute Instances. Each Compute Instance has exclusive ownership of its assigned SMs of the GPU Instance. However, all Compute Instances within a GPU Instance share the GPU Instance’s memory and memory bandwidth. Every Compute Instance acts and operates as a CUDA device with a unique device ID. See the driver release notes as well as the documentation for the <code class="docutils literal notranslate"><span class="pre">nvidia-smi</span></code> CLI tool for more information on how to configure MIG instances.</p>
<p>For profiling, a Compute Instance can be of one of two types: <em>isolated</em> or <em>shared</em>.</p>
<p>An <em>isolated</em> Compute Instance owns all of its assigned resources and does not share any GPU unit with another Compute Instance. In other words, the Compute Instance is the same size as its parent GPU Instance and consequently does not have any other sibling Compute Instances. Profiling works as usual for isolated Compute Instances.</p>
<p>A <em>shared</em> Compute Instance uses GPU resources that can potentially also be accessed by other Compute Instances in the same GPU Instance. Due to this resource sharing, collecting profiling data from those shared units is not permitted. Attempts to collect metrics from a shared unit fail with an error message of <code class="docutils literal notranslate"><span class="pre">==ERROR==</span> <span class="pre">Failed</span> <span class="pre">to</span> <span class="pre">access</span> <span class="pre">the</span> <span class="pre">following</span> <span class="pre">metrics.</span> <span class="pre">When</span> <span class="pre">profiling</span> <span class="pre">on</span> <span class="pre">a</span> <span class="pre">MIG</span> <span class="pre">instance,</span> <span class="pre">it</span> <span class="pre">is</span> <span class="pre">not</span> <span class="pre">possible</span> <span class="pre">to</span> <span class="pre">collect</span> <span class="pre">metrics</span> <span class="pre">from</span> <span class="pre">GPU</span> <span class="pre">units</span> <span class="pre">that</span> <span class="pre">are</span> <span class="pre">shared</span> <span class="pre">with</span> <span class="pre">other</span> <span class="pre">MIG</span> <span class="pre">instances</span></code> followed by the list of failing metrics. Collecting only metrics from GPU units that are exclusively owned by a shared Compute Instance is still possible.</p>
<p class="title sectiontitle rubric" id="locking-clocks">Locking Clocks</p>
<p>NVIDIA Nsight Compute is not able to set the clock frequency on any Compute Instance for profiling. You can continue analyzing kernels without fixed clock frequencies (using <code class="docutils literal notranslate"><span class="pre">--clock-control</span> <span class="pre">none</span></code>; see <a class="reference external" href="index.html#clock-control">here</a> for more details). If you have sufficient permissions, <code class="docutils literal notranslate"><span class="pre">nvidia-smi</span></code> can be used to configure a fixed frequency for the whole GPU by calling <code class="docutils literal notranslate"><span class="pre">nvidia-smi</span> <span class="pre">--lock-gpu-clocks=tdp,tdp</span></code>. This sets the GPU clocks to the base TDP frequency until you reset the clocks by calling <code class="docutils literal notranslate"><span class="pre">nvidia-smi</span> <span class="pre">--reset-gpu-clocks</span></code>.</p>
<p class="title sectiontitle rubric" id="mig-on-baremetal-non-vgpu">MIG on Baremetal (non-vGPU)</p>
<p>All Compute Instances on a GPU share the same clock frequencies.</p>
<p class="title sectiontitle rubric" id="mig-on-nvidia-vgpu">MIG on NVIDIA vGPU</p>
<p>Enabling profiling for a VM gives the VM access to the GPU’s global performance counters, which may include activity from other VMs executing on the same GPU. Enabling profiling for a VM also allows the VM to lock clocks on the GPU, which impacts all other VMs executing on the same GPU, including MIG Compute Instances.</p>
</section>
</section>
<section id="roofline-charts">
<span id="roofline"></span><h2><span class="section-number">2.8. </span>Roofline Charts<a class="headerlink" href="#roofline-charts" title="Permalink to this headline"></a></h2>
<p>Roofline charts provide a very helpful way to visualize achieved performance on complex processing units, like GPUs. This section introduces the Roofline charts that are presented within a profile report.</p>
<section id="roofline-overview">
<span id="id6"></span><h3><span class="section-number">2.8.1. </span>Overview<a class="headerlink" href="#roofline-overview" title="Permalink to this headline"></a></h3>
<p>Kernel performance is not only dependent on the operational speed of the GPU. Since a kernel requires data to work on, performance is also dependent on the rate at which the GPU can feed data to the kernel. A typical roofline chart combines the peak performance and memory bandwidth of the GPU, with a metric called <em>Arithmetic Intensity</em> (a ratio between <em>Work</em> and <em>Memory Traffic</em>), into a single chart, to more realistically represent the achieved performance of the profiled kernel. A simple roofline chart might look like the following:</p>
<figure class="align-center" id="id38">
<img alt="../_images/roofline-overview.png" src="../_images/roofline-overview.png" />
<figcaption>
<p><span class="caption-text">Roofline overview.</span><a class="headerlink" href="#id38" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p>This chart actually shows two different rooflines. However, the following components can be identified for each:</p>
<ul class="simple">
<li><p><strong>Vertical Axis</strong> - The vertical axis represents <em>Floating Point Operations per Second</em> (FLOPS). For GPUs this number can get quite large and so the numbers on this axis can be scaled for easier reading (as shown here). In order to better accommodate the range, this axis is rendered using a logarithmic scale.</p></li>
<li><p><strong>Horizontal Axis</strong> - The horizontal axis represents <em>Arithmetic Intensity</em>, which is the ratio between <em>Work</em> (expressed in floating point operations per second), and <em>Memory Traffic</em> (expressed in bytes per second). The resulting unit is in floating point operations per byte. This axis is also shown using a logarithmic scale.</p></li>
<li><p><strong>Memory Bandwidth Boundary</strong> - The memory bandwidth boundary is the <em>sloped</em> part of the roofline. By default, this slope is determined entirely by the memory transfer rate of the GPU but can be customized inside the <em>SpeedOfLight_RooflineChart.section</em> file if desired.</p></li>
<li><p><strong>Peak Performance Boundary</strong> - The peak performance boundary is the <em>flat</em> part of the roofline By default, this value is determined entirely by the peak performance of the GPU but can be customized inside the <em>SpeedOfLight_RooflineChart.section</em> file if desired.</p></li>
<li><p><strong>Ridge Point</strong> - The ridge point is the point at which the memory bandwidth boundary meets the peak performance boundary. This point is a useful reference when analyzing kernel performance.</p></li>
<li><p><strong>Achieved Value</strong> - The achieved value represents the performance of the profiled kernel. If baselines are being used, the roofline chart will also contain an achieved value for each baseline. The outline color of the plotted achieved value point can be used to determine from which baseline the point came.</p></li>
</ul>
</section>
<section id="analysis">
<span id="roofline-analysis"></span><h3><span class="section-number">2.8.2. </span>Analysis<a class="headerlink" href="#analysis" title="Permalink to this headline"></a></h3>
<p>The roofline chart can be very helpful in guiding performance optimization efforts for a particular kernel.</p>
<figure class="align-center" id="id39">
<img alt="../_images/roofline-analysis.png" src="../_images/roofline-analysis.png" />
<figcaption>
<p><span class="caption-text">Roofline anaysis.</span><a class="headerlink" href="#id39" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p>As shown here, the <em>ridge point</em> partitions the roofline chart into two regions. The area shaded in blue under the sloped <em>Memory Bandwidth Boundary</em> is the <em>Memory Bound</em> region, while the area shaded in green under the <em>Peak Performance Boundary</em> is the <em>Compute Bound</em> region. The region in which the <em>achieved value</em> falls, determines the current limiting factor of kernel performance.</p>
<p>The distance from the <em>achieved value</em> to the respective roofline boundary (shown in this figure as a dotted white line), represents the opportunity for performance improvement. The closer the <em>achieved value</em> is to the roofline boundary, the more optimal is its performance. An <em>achieved value</em> that lies on the <em>Memory Bandwidth Boundary</em> but is not yet at the height of the <em>ridge point</em> would indicate that any further improvements in overall FLOP/s are only possible if the <em>Arithmetic Intensity</em> is increased at the same time.</p>
<p>Using the baseline feature in combination with roofline charts, is a good way to track optimization progress over a number of kernel executions.</p>
</section>
</section>
<section id="memory-chart">
<h2><span class="section-number">2.9. </span>Memory Chart<a class="headerlink" href="#memory-chart" title="Permalink to this headline"></a></h2>
<p>The <em>Memory Chart</em> shows a graphical, logical representation of performance data for memory subunits on and off the GPU. Performance data includes transfer sizes, hit rates, number of instructions or requests, etc.</p>
<section id="memory-chart-overview">
<span id="id7"></span><h3><span class="section-number">2.9.1. </span>Overview<a class="headerlink" href="#memory-chart-overview" title="Permalink to this headline"></a></h3>
<figure class="align-center" id="id40">
<img alt="../_images/memory-chart-a100.png" src="../_images/memory-chart-a100.png" />
<figcaption>
<p><span class="caption-text">Memory chart for an NVIDIA A100 GPU</span><a class="headerlink" href="#id40" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="rubric-h3 rubric" id="logical-units-green">Logical Units (green)</p>
<p>Logical units are shown in green color.</p>
<ul class="simple">
<li><p>Kernel: The CUDA kernel executing on the GPU’s Streaming Multiprocessors</p></li>
<li><p>Global: CUDA <a class="reference external" href="index.html#metrics-hw-memory">global memory</a></p></li>
<li><p>Local: CUDA <a class="reference external" href="index.html#metrics-hw-memory">local memory</a></p></li>
<li><p>Texture: CUDA <a class="reference external" href="index.html#metrics-hw-tex-surf">texture memory</a></p></li>
<li><p>Surface: CUDA <a class="reference external" href="index.html#metrics-hw-tex-surf">surface memory</a></p></li>
<li><p>Shared: CUDA <a class="reference external" href="index.html#metrics-hw-memory">shared memory</a></p></li>
<li><p>Load Global Store Shared: Instructions loading directly from global into shared memory without intermediate register file access</p></li>
</ul>
<p class="rubric-h3 rubric" id="physical-units-blue">Physical Units (blue)</p>
<p>Physical units are shown in blue color.</p>
<ul class="simple">
<li><p>L1/TEX Cache: The <a class="reference external" href="index.html#metrics-hw-caches">L1/Texture cache</a>. The underlying physical memory is split between this cache and the user-managed <em>Shared Memory</em>.</p></li>
<li><p>Shared Memory: CUDA’s user-managed <a class="reference external" href="index.html#metrics-hw-memory">shared memory</a>. The underlying physical memory is split between this and the <em>L1/TEX Cache</em>.</p></li>
<li><p>L2 Cache: The <a class="reference external" href="index.html#metrics-hw-caches">L2 cache</a></p></li>
<li><p>L2 Compression: The memory compression unit of the <em>L2 Cache</em></p></li>
<li><p>System Memory: Off-chip <a class="reference external" href="index.html#metrics-hw-memory">system (CPU) memory</a></p></li>
<li><p>Device Memory: On-chip <a class="reference external" href="index.html#metrics-hw-memory">device (GPU) memory</a> of the CUDA device that executes the kernel</p></li>
<li><p>Peer Memory: On-chip <a class="reference external" href="index.html#metrics-hw-memory">device (GPU) memory</a> of other CUDA devices</p></li>
</ul>
<p>Depending on the exact GPU architecture, the exact set of shown units can vary, as not all GPUs have all units.</p>
<p class="rubric-h3 rubric" id="links">Links</p>
<p>Links between <em>Kernel</em> and other logical units represent the number of executed instructions (<em>Inst</em>) targeting the respective unit. For example, the link between <em>Kernel</em> and <em>Global</em> represents the instructions loading from or storing to the global memory space. Instructions using the NVIDIA A100’s <em>Load Global Store Shared</em> paradigm are shown separately, as their register or cache access behavior can be different from regular global loads or shared stores.</p>
<p>Links between logical units and blue, physical units represent the number of requests (<em>Req</em>) issued as a result of their respective instructions. For example, the link going from <em>L1/TEX Cache</em> to <em>Global</em> shows the number of requests generated due to global load instructions.</p>
<p>The color of each link represents the percentage of peak utilization of the corresponding communication path. The color legend to the right of the chart shows the applied color gradient from unused (0%) to operating at peak performance (100%). Triangle markers to the left of the legend correspond to the links in the chart. The markers offer a more accurate value estimate for the achieved peak performances than the color gradient alone.</p>
<p>A unit often shares a common data port for incoming and outgoing traffic. While the links sharing a port might operate well below their individual peak performances, the unit’s data port may have already reached its peak. Port utilization is shown in the chart by colored rectangles inside the units located at the incoming and outgoing links. Ports use the same color gradient as the data links and have also a corresponding marker to the left of the legend.</p>
<p>An example of the correlation between the peak values reported in the memory tables and the ports in the memory chart is shown below.</p>
<figure class="align-center" id="id41">
<img alt="../_images/memory-peak-mapping.png" src="../_images/memory-peak-mapping.png" />
<figcaption>
<p><span class="caption-text">Mapping of peak values between memory tables and memory chart</span><a class="headerlink" href="#id41" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="rubric-h2 rubric" id="metrics">Metrics</p>
<p>Metrics from this chart can be collected on the command line using <code class="docutils literal notranslate"><span class="pre">--set</span> <span class="pre">full</span></code>, <code class="docutils literal notranslate"><span class="pre">--section</span> <span class="pre">MemoryWorkloadAnalysis_Chart</span></code> or <code class="docutils literal notranslate"><span class="pre">--metrics</span> <span class="pre">group:memory__chart</span></code>.</p>
</section>
</section>
<section id="memory-tables">
<h2><span class="section-number">2.10. </span>Memory Tables<a class="headerlink" href="#memory-tables" title="Permalink to this headline"></a></h2>
<p>The <em>Memory Tables</em> show detailed metrics for the various memory HW units, such as shared memory, the caches, and device memory. For most table entries, you can hover over it to see the underlying metric name and description. Some entries are generated as derivatives from other cells, and do not show a metric name on their own, but the respective calculation. If a certain metric does not contribute to the generic derivative calculation, it is shown as <em>UNUSED</em> in the tooltip. You can hover over row or column headers to see a description of this part of the table.</p>
<section id="shared-memory">
<span id="memory-tables-smem"></span><h3><span class="section-number">2.10.1. </span>Shared Memory<a class="headerlink" href="#shared-memory" title="Permalink to this headline"></a></h3>
<figure class="align-center" id="id42">
<img alt="../_images/memory-tables-smem.png" src="../_images/memory-tables-smem.png" />
<figcaption>
<p><span class="caption-text">Example Shared Memory table, collected on an RTX 2080 Ti</span><a class="headerlink" href="#id42" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="rubric-h3 rubric" id="columns">Columns</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 10%" />
<col style="width: 90%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Instructions</span></code></p></td>
<td><p>For each access type, the total number of all actually executed assembly (SASS) <a class="reference external" href="index.html#metrics-quantities">instructions</a> per warp. Predicated-off instructions are not included.</p>
<p>E.g., the instruction <em>STS</em> would be counted towards <em>Shared Store</em>.</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Requests</span></code></p></td>
<td><p>The total number of all <a class="reference external" href="index.html#metrics-quantities">requests</a> to shared memory. On SM 7.0 (Volta) and newer architectures, each shared memory instruction generates exactly one request.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Wavefronts</span></code></p></td>
<td><p>Number of <a class="reference external" href="index.html#metrics-quantities">wavefronts</a> required to service the requested shared memory data. Wavefronts are serialized and processed on different cycles.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">%</span> <span class="pre">Peak</span></code></p></td>
<td><p>Percentage of peak utilization. Higher values imply a higher utilization of the unit and can show potential bottlenecks, as it does not necessarily indicate efficient usage.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Bank</span> <span class="pre">Conflicts</span></code></p></td>
<td><p>If multiple threads’ requested addresses map to different offsets in the same memory bank, the accesses are serialized. The hardware splits a conflicting memory request into as many separate conflict-free requests as necessary, decreasing the effective bandwidth by a factor equal to the number of colliding memory requests.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="rows">Rows</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 26%" />
<col style="width: 74%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">(Access</span> <span class="pre">Types)</span></code></p></td>
<td><p>Shared memory access operations.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Total</span></code></p></td>
<td><p>The aggregate for all access types in the same column.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-1">Metrics</p>
<p>Metrics from this table can be collected on the command line using <code class="docutils literal notranslate"><span class="pre">--set</span> <span class="pre">full</span></code>, <code class="docutils literal notranslate"><span class="pre">--section</span> <span class="pre">MemoryWorkloadAnalysis_Tables</span></code> or <code class="docutils literal notranslate"><span class="pre">--metrics</span> <span class="pre">group:memory__shared_table</span></code>.</p>
</section>
<section id="l1-tex-cache">
<span id="memory-tables-l1"></span><h3><span class="section-number">2.10.2. </span>L1/TEX Cache<a class="headerlink" href="#l1-tex-cache" title="Permalink to this headline"></a></h3>
<figure class="align-center" id="id43">
<img alt="../_images/memory-tables-l1.png" src="../_images/memory-tables-l1.png" />
<figcaption>
<p><span class="caption-text">Example L1/TEX Cache memory table, collected on an RTX 2080 Ti</span><a class="headerlink" href="#id43" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<figure class="align-center" id="id44">
<img alt="../_images/hw-model-l1tex-ga100-global.png" src="../_images/hw-model-l1tex-ga100-global.png" />
<figcaption>
<p><span class="caption-text">Model of the Global Load Pipeline for the L1TEX cache on GA100, mapped to the memory table.</span><a class="headerlink" href="#id44" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="rubric-h3 rubric" id="columns-1">Columns</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 6%" />
<col style="width: 94%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Instructions</span></code></p></td>
<td><p>For each access type, the total number of all actually executed assembly (SASS) <a class="reference external" href="index.html#metrics-quantities">instructions</a> per warp. Predicated-off instructions are not included.</p>
<p>E.g., the instruction <em>LDG</em> would be counted towards <em>Global Loads</em>.</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Requests</span></code></p></td>
<td><p>The total number of all <a class="reference external" href="index.html#metrics-quantities">requests</a> to L1, generated for each instruction type. On SM 7.0 (Volta) and newer architectures, each instruction generates exactly one request for LSU traffic (global, local, …). For texture (TEX) traffic, more than one request may be generated.</p>
<p>In the example, each of the 65536 global load instructions generates exactly one request.</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Wavefronts</span></code></p></td>
<td><p>Number of <a class="reference external" href="index.html#metrics-quantities">wavefronts</a> required to service the requested memory operation. Wavefronts are serialized and processed on different cycles.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Wavefront</span> <span class="pre">%</span> <span class="pre">Peak</span></code></p></td>
<td><p>Percentage of peak utilization for the units processing <a class="reference external" href="index.html#metrics-quantities">wavefronts</a>. High numbers can imply that the processing pipelines are saturated and can become a bottleneck.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Sectors</span></code></p></td>
<td><p>The total number of all L1 <a class="reference external" href="index.html#metrics-quantities">sectors</a> accesses sent to L1. Each load or store request accesses one or more sectors in the L1 cache. Atomics and reductions are passed through to the L2 cache.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Sectors/Req</span></code></p></td>
<td><p>The average ratio of sectors to requests for the L1 cache. For the same number of active threads in a warp, smaller numbers imply a more efficient memory access pattern. For warps with 32 active threads, the optimal ratios per access size are: 32-bit: 4, 64-bit: 8, 128-bit: 16. Smaller ratios indicate some degree of uniformity or overlapped loads within a cache line. Higher numbers can imply <a class="reference external" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory-accesses">uncoalesced memory accesses</a> and will result in increased memory traffic.</p>
<p>In the example, the average ratio for global loads is 32 sectors per request, which implies that each thread needs to access a different sector. Ideally, for warps with 32 active threads, with each thread accessing a single, aligned 32-bit value, the ratio would be 4, as every 8 consecutive threads access the same sector.</p>
</td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Hit</span> <span class="pre">Rate</span></code></p></td>
<td><p><a class="reference external" href="index.html#metrics-quantities">Sector</a> hit rate (percentage of requested sectors that do not miss) in the L1 cache. Sectors that miss need to be requested from L2, thereby contributing to <em>Sector Misses to L2</em>. Higher hit rates imply better performance due to lower access latencies, as the request can be served by L1 instead of a later stage. Not to be confused with <em>Tag Hit Rate</em> (not shown).</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Bytes</span></code></p></td>
<td><p>Total number of bytes requested from L1. This is identical to the number of sectors multiplied by 32 byte, since the minimum access size in L1 is one sector.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Sector</span> <span class="pre">Misses</span> <span class="pre">to</span> <span class="pre">L2</span></code></p></td>
<td><p>Total number of sectors that miss in L1 and generate subsequent requests in the <a class="reference external" href="index.html#memory-tables-l2">L2 Cache</a>.</p>
<p>In this example, the 262144 sector misses for global and local loads can be computed as the miss-rate of 12.5%, multiplied by the number of 2097152 sectors.</p>
</td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">%</span> <span class="pre">Peak</span> <span class="pre">to</span> <span class="pre">L2</span></code></p></td>
<td><p>Percentage of peak utilization of the L1-to-XBAR interface, used to send L2 cache requests. If this number is high, the workload is likely dominated by scattered {writes, atomics, reductions}, which can increase the latency and cause <a class="reference external" href="index.html#statistical-sampler__warp-scheduler-states">warp stalls</a>.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Returns</span> <span class="pre">to</span> <span class="pre">SM</span></code></p></td>
<td><p>Number of return packets sent from the L1 cache back to the SM. Larger request access sizes result in higher number of returned packets.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">%</span> <span class="pre">Peak</span> <span class="pre">to</span> <span class="pre">SM</span></code></p></td>
<td><p>Percentage of peak utilization of the XBAR-to-L1 return path (compare Returns to SM). If this number is high, the workload is likely dominated by scattered reads, thereby causing <a class="reference external" href="index.html#statistical-sampler__warp-scheduler-states">warp stalls</a>. Improving read-coalescing or the <em>L1 hit rate</em> could reduce this utilization.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="rows-1">Rows</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 17%" />
<col style="width: 83%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">(Access</span> <span class="pre">Types)</span></code></p></td>
<td><p>The various access types, e.g. loads from global memory or reduction operations on surface memory.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Loads</span></code></p></td>
<td><p>The aggregate of all load access types in the same column.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Stores</span></code></p></td>
<td><p>The aggregate of all store access types in the same column.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Total</span></code></p></td>
<td><p>The aggregate of all load and store access types in the same column.</p></td>
</tr>
</tbody>
</table>
<p class="rubric-h3 rubric" id="metrics-2">Metrics</p>
<p>Metrics from this table can be collected on the command line using <code class="docutils literal notranslate"><span class="pre">--set</span> <span class="pre">full</span></code>, <code class="docutils literal notranslate"><span class="pre">--section</span> <span class="pre">MemoryWorkloadAnalysis_Tables</span></code> or <code class="docutils literal notranslate"><span class="pre">--metrics</span> <span class="pre">group:memory__first_level_cache_table</span></code>.</p>
</section>
<section id="l2-cache">
<span id="memory-tables-l2"></span><h3><span class="section-number">2.10.3. </span>L2 Cache<a class="headerlink" href="#l2-cache" title="Permalink to this headline"></a></h3>
<figure class="align-center" id="id45">
<img alt="../_images/memory-tables-l2.png" src="../_images/memory-tables-l2.png" />
<figcaption>
<p><span class="caption-text">Example L2 Cache memory table, collected on an RTX 2080 Ti</span><a class="headerlink" href="#id45" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<figure class="align-center" id="id46">
<img alt="../_images/hw-model-lts-ga100.png" src="../_images/hw-model-lts-ga100.png" />
<figcaption>
<p><span class="caption-text">Model of the L2 cache on GA100, mapped to the memory table.</span><a class="headerlink" href="#id46" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="title sectiontitle rubric" id="columns-2">Columns</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 5%" />
<col style="width: 95%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Requests</span></code></p></td>
<td><p>For each access type, the total number of <a class="reference external" href="index.html#metrics-decoder__metrics-quantities">requests</a> made to the L2 cache. This correlates with the <a class="reference external" href="index.html#memory-tables-l1__memory-tables-l1-columns">Sector Misses to L2</a> for the L1 cache. Each request accesses up to four sectors from a single 128 byte cache line.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Sectors</span></code></p></td>
<td><p>For each access type, the total number of <a class="reference external" href="index.html#metrics-decoder__metrics-quantities">sectors</a> requested from the L2 cache. Each request accesses between one and four sectors.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Sectors/Req</span></code></p></td>
<td><p>The average ratio of sectors to requests for the L2 cache. For the same number of active threads in a warp, smaller numbers imply a more efficient memory access pattern. Smaller ratios indicate some degree of uniformity or overlapped loads within a cache line. Higher numbers can imply <a class="reference external" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory-accesses">uncoalesced memory accesses</a> and will result in increased memory traffic.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">%</span> <span class="pre">Peak</span></code></p></td>
<td><p>Percentage of peak sustained number of sectors. The “work package” in the L2 cache is a sector. Higher values imply a higher utilization of the unit and can show potential bottlenecks, as it does not necessarily indicate efficient usage.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Hit</span> <span class="pre">Rate</span></code></p></td>
<td><p>Hit rate (percentage of requested sectors that do not miss) in the L2 cache. Sectors that miss need to be requested from a later stage, thereby contributing to one of <em>Sector Misses to Device</em>, <em>Sector Misses to System</em>, or <em>Sector Misses to Peer</em>. Higher hit rates imply better performance due to lower access latencies, as the request can be served by L2 instead of a later stage.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Bytes</span></code></p></td>
<td><p>Total number of bytes requested from L2. This is identical to the number of sectors multiplied by 32 byte, since the minimum access size in L2 is one sector.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Throughput</span></code></p></td>
<td><p>Achieved L2 cache throughput in bytes per second. High values indicate high utilization of the unit.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Sector</span> <span class="pre">Misses</span> <span class="pre">to</span> <span class="pre">Device</span></code></p></td>
<td><p>Total number of sectors that miss in L2 and generate <a class="reference external" href="index.html#memory-tables-dram">subsequent requests</a> in <a class="reference external" href="index.html#metrics-hw-model__metrics-hw-memory">device memory</a>.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Sector</span> <span class="pre">Misses</span> <span class="pre">to</span> <span class="pre">System</span></code></p></td>
<td><p>Total number of sectors that miss in L2 and generate subsequent requests in <a class="reference external" href="index.html#metrics-hw-model__metrics-hw-memory">system memory</a>.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Sector</span> <span class="pre">Misses</span> <span class="pre">to</span> <span class="pre">Peer</span></code></p></td>
<td><p>Total number of sectors that miss in L2 and generate subsequent requests in <a class="reference external" href="index.html#metrics-hw-model__metrics-hw-memory">peer memory</a>.</p></td>
</tr>
</tbody>
</table>
<p class="title sectiontitle rubric" id="rows-2">Rows</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 8%" />
<col style="width: 92%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">(Access</span> <span class="pre">Types)</span></code></p></td>
<td><p>The various access types, e.g. loads or reductions originating from L1 cache.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">L1/TEX</span> <span class="pre">Total</span></code></p></td>
<td><p>Total for all operations originating from the L1 cache.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">ECC</span> <span class="pre">Total</span></code></p></td>
<td><p>Total for all operations caused by ECC (Error Correction Code). If ECC is enabled, L2 write requests that partially modify a sector cause a corresponding sector load from DRAM. These additional load operations increase the sector misses of L2.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">L2</span> <span class="pre">Fabric</span> <span class="pre">Total</span></code></p></td>
<td><p>Total for all operations across the L2 fabric connecting the two L2 partitions. This row is only shown for kernel launches on CUDA devices with L2 fabric.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">GPU</span> <span class="pre">Total</span></code></p></td>
<td><p>Total for all operations across all clients of the L2 cache. Independent of having them split out separately in this table.</p></td>
</tr>
</tbody>
</table>
<p class="title sectiontitle rubric" id="metrics-3">Metrics</p>
<p>Metrics from this table can be collected on the command line using <code class="docutils literal notranslate"><span class="pre">--set</span> <span class="pre">full</span></code>, <code class="docutils literal notranslate"><span class="pre">--section</span> <span class="pre">MemoryWorkloadAnalysis_Tables</span></code> or <code class="docutils literal notranslate"><span class="pre">--metrics</span> <span class="pre">group:memory__l2_cache_table</span></code>.</p>
</section>
<section id="l2-cache-eviction-policies">
<span id="memory-tables-l2-evict-policy"></span><h3><span class="section-number">2.10.4. </span>L2 Cache Eviction Policies<a class="headerlink" href="#l2-cache-eviction-policies" title="Permalink to this headline"></a></h3>
<figure class="align-center" id="id47">
<img alt="../_images/memory-tables-l2-evict-policy.png" src="../_images/memory-tables-l2-evict-policy.png" />
<figcaption>
<p><span class="caption-text">Example L2 Cache Eviction Policies memory table, collected on an A100 GPU</span><a class="headerlink" href="#id47" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="title sectiontitle rubric" id="columns-3">Columns</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 5%" />
<col style="width: 95%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">First</span></code></p></td>
<td><p>Number of sectors accessed in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_first</span></code> policy. Data cached with this policy will be first in the eviction priority order and will likely be evicted when cache eviction is required. This policy is suitable for streaming data.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Hit</span> <span class="pre">Rate</span></code></p></td>
<td><p>Cache hit rate for sector accesses in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_first</span></code> policy.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Last</span></code></p></td>
<td><p>Number of sectors accessed in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_last</span></code> policy. Data cached with this policy will be last in the eviction priority order and will likely be evicted only after other data with <code class="docutils literal notranslate"><span class="pre">evict_normal</span></code> or <code class="docutils literal notranslate"><span class="pre">evict_first</span></code> eviction policy is already evicted. This policy is suitable for data that should remain persistent in cache.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Hit</span> <span class="pre">Rate</span></code></p></td>
<td><p>Cache hit rate for sector accesses in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_last</span></code> policy.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Normal</span></code></p></td>
<td><p>Number of sectors accessed in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_normal</span></code> policy. This is the default policy.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Hit</span> <span class="pre">Rate</span></code></p></td>
<td><p>Cache hit rate for sector accesses in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_normal</span></code> policy.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Normal</span> <span class="pre">Demote</span></code></p></td>
<td><p>Number of sectors accessed in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_normal_demote</span></code> policy.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Hit</span> <span class="pre">Rate</span></code></p></td>
<td><p>Cache hit rate for sector accesses in the L2 cache using the <code class="docutils literal notranslate"><span class="pre">evict_normal_demote</span></code> policy.</p></td>
</tr>
</tbody>
</table>
<p class="title sectiontitle rubric" id="rows-3">Rows</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 12%" />
<col style="width: 88%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">(Access</span> <span class="pre">Types)</span></code></p></td>
<td><p>The various access types, e.g. loads or reductions, originating from L1 cache.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">L1/TEX</span> <span class="pre">Total</span></code></p></td>
<td><p>Total for all operations originating from the L1 cache.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">L2</span> <span class="pre">Fabric</span> <span class="pre">Total</span></code></p></td>
<td><p>Total for all operations across the L2 fabric connecting the two L2 partitions. This row is only shown for kernel launches on CUDA devices with L2 fabric.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">GPU</span> <span class="pre">Total</span></code></p></td>
<td><p>Total for all operations across all clients of the L2 cache. Independent of having them split out separately in this table.</p></td>
</tr>
</tbody>
</table>
<p class="title sectiontitle rubric" id="metrics-4">Metrics</p>
<p>Metrics from this table can be collected on the command line using <code class="docutils literal notranslate"><span class="pre">--set</span> <span class="pre">full</span></code>, <code class="docutils literal notranslate"><span class="pre">--section</span> <span class="pre">MemoryWorkloadAnalysis_Tables</span></code> or <code class="docutils literal notranslate"><span class="pre">--metrics</span> <span class="pre">group:memory__l2_cache_evict_policy_table</span></code>. Note that this table is only available on GPUs with GA100 or newer.</p>
</section>
<section id="device-memory">
<h3><span class="section-number">2.10.5. </span>Device Memory<a class="headerlink" href="#device-memory" title="Permalink to this headline"></a></h3>
<figure class="align-center" id="id48">
<img alt="../_images/memory-tables-dram.png" src="../_images/memory-tables-dram.png" />
<figcaption>
<p><span class="caption-text">Example Device Memory table, collected on an RTX 2080 Ti</span><a class="headerlink" href="#id48" title="Permalink to this image"></a></p>
</figcaption>
</figure>
<p class="title sectiontitle rubric" id="columns-4">Columns</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 8%" />
<col style="width: 92%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Sectors</span></code></p></td>
<td><p>For each access type, the total number of <a class="reference external" href="index.html#metrics-decoder__metrics-quantities">sectors</a> requested from device memory.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">%</span> <span class="pre">Peak</span></code></p></td>
<td><p>Percentage of peak device memory utilization. Higher values imply a higher utilization of the unit and can show potential bottlenecks, as it does not necessarily indicate efficient usage.</p></td>
</tr>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">Bytes</span></code></p></td>
<td><p>Total number of bytes transferred between <a class="reference external" href="index.html#memory-tables-l2">L2 Cache</a> and device memory.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Throughput</span></code></p></td>
<td><p>Achieved device memory throughput in bytes per second. High values indicate high utilization of the unit.</p></td>
</tr>
</tbody>
</table>
<p class="title sectiontitle rubric" id="rows-4">Rows</p>
<table class="table-no-stripes docutils align-default">
<colgroup>
<col style="width: 26%" />
<col style="width: 74%" />
</colgroup>
<tbody>
<tr class="row-odd"><td><p><code class="docutils literal notranslate"><span class="pre">(Access</span> <span class="pre">Types)</span></code></p></td>
<td><p>Device memory loads and stores.</p></td>
</tr>
<tr class="row-even"><td><p><code class="docutils literal notranslate"><span class="pre">Total</span></code></p></td>
<td><p>The aggregate for all access types in the same column.</p></td>
</tr>
</tbody>
</table>
<p class="title sectiontitle rubric" id="metrics-5">Metrics</p>
<p>Metrics from this table can be collected on the command line using <code class="docutils literal notranslate"><span class="pre">--set</span> <span class="pre">full</span></code>, <code class="docutils literal notranslate"><span class="pre">--section</span> <span class="pre">MemoryWorkloadAnalysis_Tables</span></code> or <code class="docutils literal notranslate"><span class="pre">--metrics</span> <span class="pre">group:memory__dram_table</span></code>.</p>
</section>
</section>
<section id="faq">
<h2><span class="section-number">2.11. </span>FAQ<a class="headerlink" href="#faq" title="Permalink to this headline"></a></h2>
<ul>
<li><p><strong>n/a metric values</strong></p>
<p>n/a means that the metric value is “not available”. The most common reason is that the requested metric does not exist. This can either be the result of a typo, or a missing <a class="reference external" href="index.html#metrics-structure__metrics-examples">suffix</a>. Verify the metric name against the output of of the <code class="docutils literal notranslate"><span class="pre">--query-metrics</span></code>NVIDIA Nsight Compute CLI option.</p>
<p>If the metric name was copied (e.g. from an old version of this documentation), make sure that it does not contain zero-width unicode characters.</p>
<p>Finally, the metric might simply not exist for the targeted GPU architecture. For example, the IMMA pipeline metric <code class="docutils literal notranslate"><span class="pre">sm__inst_executed_pipe_tensor_op_imma.avg.pct_of_peak_sustained_active</span></code> is not available on GV100 chips.</p>
</li>
<li><p><strong>Metric values outside the expected logical range</strong></p>
<p>This includes e.g. percentages exceeding 100% or metrics reporting negative values. For further details, see <a class="reference external" href="index.html#range-and-precision">Range and Precision</a>.</p>
</li>
<li><p><strong>ERR_NVGPUCTRPERM - The user does not have permission to access NVIDIA GPU Performance Counters on the target device.</strong></p>
<p>By default, NVIDIA drivers require elevated permissions to access GPU performance counters. On mobile platforms, profile as root/using sudo. On other platforms, you can either start profiling as root/using sudo, or by enabling non-admin profiling. For further details, see <a class="reference external" href="https://developer.nvidia.com/ERR_NVGPUCTRPERM">https://developer.nvidia.com/ERR_NVGPUCTRPERM</a>.</p>
<p>On Windows Subsystem for Linux (WSL), access to NVIDIA GPU Performance Counters must be enabled in the NVIDIA Control Panel of the Windows host.</p>
</li>
<li><p><strong>Unsupported GPU</strong></p>
<p>This indicates that the GPU, on which the current kernel is launched, is not supported. See the <a class="reference external" href="../ReleaseNotes/index.html#gpu-support">Release Notes</a> for a list of devices supported by your version of NVIDIA Nsight Compute. It can also indicate that the current <em>GPU configuration</em> is not supported. For example, NVIDIA Nsight Compute might not be able to profile GPUs in SLI configuration.</p>
</li>
<li><p><strong>Connection error detected communicating with target application.</strong></p>
<p>The inter-process connection to the profiled application unexpectedly dropped. This happens if the application is killed or signals an exception (e.g. segmentation fault).</p>
</li>
<li><p><strong>Failed to connect. The target process may have exited.</strong></p>
<p>This occurs if</p>
<ul class="simple">
<li><p>the application does not call any CUDA API calls before it exits.</p></li>
<li><p>the application terminates early because it was started from the wrong working directory, or with the wrong arguments. In this case, check the details in the <em>Connection Dialog</em>.</p></li>
<li><p>the application crashes before calling any CUDA API calls.</p></li>
<li><p>the application launches child processes which use the CUDA. In this case, launch with the <code class="docutils literal notranslate"><span class="pre">--target-processes</span> <span class="pre">all</span></code> option.</p></li>
</ul>
</li>
<li><p><strong>The profiler returned an error code: (number)</strong></p>
<p>For the non-interactive <em>Profile</em> activity, the NVIDIA Nsight Compute CLI is started to generate the report. If either the application exited with a non-zero return code, or the NVIDIA Nsight Compute CLI encountered an error itself, the resulting return code will be shown in this message.</p>
<p>For example, if the application hit a segmentation fault (SIGSEGV) on Linux, it will likely return error code 11. All non-zero return codes are considered errors, so the message is also shown if the application exits with return code 1 during regular execution.</p>
<p>To debug this issue, it can help to run the data collection directly from the command line using <code class="docutils literal notranslate"><span class="pre">ncu</span></code> in order to observe the application’s and the profiler’s command line output, e.g. <code class="docutils literal notranslate"><span class="pre">==ERROR==</span> <span class="pre">The</span> <span class="pre">application</span> <span class="pre">returned</span> <span class="pre">an</span> <span class="pre">error</span> <span class="pre">code</span> <span class="pre">(11)</span></code></p>
</li>
<li><p><strong>Failed to open/create lock file (path). Please check that this process has write permissions on this file.</strong></p>
<p>NVIDIA Nsight Compute failed to create or open the file <code class="docutils literal notranslate"><span class="pre">(path)</span></code> with write permissions. This file is used for inter-process <a class="reference external" href="index.html#serialization">serialization</a>. NVIDIA Nsight Compute does not remove this file after profiling by design. The error occurs if the file was created by a profiling process with permissions that prevent the current process from writing to this file, or if the current user can’t acquire this file for other reasons (e.g. certain Linux kernel security settings).</p>
<p>The file is in the current temporary directory, i.e. <code class="docutils literal notranslate"><span class="pre">TMPDIR/nsight-compute-lock</span></code>. On Windows, <em>TMPDIR</em> is the path returned by the Windows <code class="docutils literal notranslate"><span class="pre">GetTempPath</span></code> API function. On other platforms, it is the path supplied by the first environment variable in the list <code class="docutils literal notranslate"><span class="pre">TMPDIR,</span> <span class="pre">TMP,</span> <span class="pre">TEMP,</span> <span class="pre">TEMPDIR</span></code>. If none of these is found, it’s <code class="docutils literal notranslate"><span class="pre">/var/nvidia</span></code> on QNX and <code class="docutils literal notranslate"><span class="pre">/tmp</span></code> otherwise.</p>
<p>Older versions of NVIDIA Nsight Compute did not set write permissions for all users on this file by default. As a result, running the tool on the same system with a different user might cause this error. This has been resolved since version 2020.2.1.</p>
<p>The following workarounds can be used to solve this problem:</p>
<ul class="simple">
<li><p>If it is otherwise ensured that no concurrent NVIDIA Nsight Compute instances are active on the same system, set <code class="docutils literal notranslate"><span class="pre">TMPDIR</span></code> to a different directory for which the current user has write permissions.</p></li>
<li><p>Ask the user owning the file, or a system administrator, to remove it or add write permissions for all potential users.</p></li>
<li><p>On Linux systems setting <code class="docutils literal notranslate"><span class="pre">fs.protected_regular=1</span></code>, root or other users may not be able to access this file, even though the owner can, if the sticky bit is set on the temporary directory. Either disable this setting using <code class="docutils literal notranslate"><span class="pre">sudo</span> <span class="pre">sysctl</span> <span class="pre">fs.protected_regular=0</span></code>, use a different temporary directory (see above), or enable access to hardware performance counters for non-root users and profile as the same user who owns the file (see <a class="reference external" href="https://developer.nvidia.com/ERR_NVGPUCTRPERM">https://developer.nvidia.com/ERR_NVGPUCTRPERM</a> on how to change this setting).</p></li>
</ul>
</li>
<li><p><strong>Profiling failed because a driver resource was unavailable.</strong></p>
<p>The error indicates that a required CUDA driver resource was unavailable during profiling. Most commonly, this means that NVIDIA Nsight Compute could not reserve the driver’s performance monitor, which is necessary for collecting most metrics.</p>
<p>This can happen if another application has a concurrent reservation on this resource. Such applications can be e.g. <a class="reference external" href="https://developer.nvidia.com/dcgm">DCGM</a>, a client of <a class="reference external" href="https://developer.nvidia.com/cupti">CUPTI’s Profiling API</a>, <a class="reference external" href="https://developer.nvidia.com/nsight-graphics">Nsight Graphics</a>, or another instance of NVIDIA Nsight Compute without access to the same file system (see <a class="reference external" href="index.html#serialization">serialization</a> for how this is prevented within the same file system).</p>
<p>If you expect the problem to be caused by DCGM, consider using <code class="docutils literal notranslate"><span class="pre">dcgmi</span> <span class="pre">profile</span> <span class="pre">--pause</span></code> to stop its monitoring while profiling with NVIDIA Nsight Compute.</p>
</li>
<li><p><strong>Could not deploy stock * files to *</strong></p>
<p><strong>Could not determine user home directory for section deployment.</strong></p>
<p>An error occurred while trying to deploy stock section or rule files. By default, NVIDIA Nsight Compute tries to deploy these to a versioned directory in the user’s home directory (as identified by the <code class="docutils literal notranslate"><span class="pre">HOME</span></code> environment variable on Linux), e.g. <code class="docutils literal notranslate"><span class="pre">/home/user/Documents/NVIDIA</span> <span class="pre">Nsight</span> <span class="pre">Compute/<version>/Sections</span></code>.</p>
<p>If the directory cannot be determined (e.g. because this environment variable is not pointing to a valid directory), or if there is an error while deploying the files (e.g. because the current process does not have write permissions on it), warning messages are shown and NVIDIA Nsight Compute falls back to using stock sections and rules from the installation directory.</p>
<p>If you are in an environment where you consistently don’t have write access to the user’s home directory, consider populating this directory upfront using <code class="docutils literal notranslate"><span class="pre">ncu</span> <span class="pre">--section-folder-restore</span></code>, or by making <code class="docutils literal notranslate"><span class="pre">/home/user/Documents/NVIDIA</span> <span class="pre">Nsight</span> <span class="pre">Compute/<version></span></code> a symlink to a writable directory.</p>
</li>
<li><p><strong>ProxyJump SSH option is not working</strong></p>
<p>NVIDIA Nsight Compute does not manage authentication or interactive prompts with the OpenSSH client launched when using the <code class="docutils literal notranslate"><span class="pre">ProxyJump</span></code> option. Therefore, to connect through an intermediate host for the first time, you will not be able to accept the intermediate host’s key. A simple way to pinpoint the cause of failures in this case is to open a terminal and use the OpenSSH client to connect to the remote target. Once that connection succeeds, NVIDIA Nsight Compute should be able to connect to the target, too.</p>
</li>
<li><p><strong>SSH connection fails without trying to connect</strong></p>
<p>If the connection fails without trying to connect, there may be a problem with the settings you entered into the connection dialog. Please make sure that the <code class="docutils literal notranslate"><span class="pre">IP/Host</span> <span class="pre">Name</span></code>, <code class="docutils literal notranslate"><span class="pre">User</span> <span class="pre">Name</span></code> and <code class="docutils literal notranslate"><span class="pre">Port</span></code> fields are correctly set.</p>
</li>
<li><p><strong>SSH connections are still not working</strong></p>
<p>The problem might come from NVIDIA Nsight Compute’s SSH client not finding a suitable host key algorithm to use which is supported by the remote server. You can force NVIDIA Nsight Compute to use a specific set of host key algorithms by setting the <code class="docutils literal notranslate"><span class="pre">HostKeyAlgorithms</span></code> option for the problematic host in your SSH configuration file. To list the supported host key algorithms for a remote target, you can use the <code class="docutils literal notranslate"><span class="pre">ssh-keyscan</span></code> utility which comes with the OpenSSH client.</p>
</li>
<li><p><strong>Removing host keys from known hosts files</strong></p>
<p>When connecting to a target machine, NVIDIA Nsight Compute tries to verify the target’s host key against the same local database as the OpenSSH client. If NVIDIA Nsight Compute find the host key is incorrect, it will inform you through a failure dialog. If you trust the key hash shown in the dialog, you can remove the previously saved key for that host by manually editing your known hosts database or using the <code class="docutils literal notranslate"><span class="pre">ssh-keygen</span> <span class="pre">-R</span> <span class="pre"><host></span></code> command.</p>
</li>
<li><p><strong>Qt initialization failed</strong></p>
<p><strong>Failed to load Qt platform plugin</strong></p>
<p>See <a class="reference external" href="../ReleaseNotes/index.html#system-requirements">System Requirements</a> for Linux.</p>
</li>
</ul>
<p class="rubric-h1 rubric">Notices</p>
<p class="rubric-h2 rubric">Notices</p>
<p>ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS, DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND SEPARATELY, “MATERIALS”) ARE BEING PROVIDED “AS IS.” NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE.</p>
<p>Information furnished is believed to be accurate and reliable. However, NVIDIA Corporation assumes no responsibility for the consequences of use of such information or for any infringement of patents or other rights of third parties that may result from its use. No license is granted by implication of otherwise under any patent rights of NVIDIA Corporation. Specifications mentioned in this publication are subject to change without notice. This publication supersedes and replaces all other information previously supplied. NVIDIA Corporation products are not authorized as critical components in life support devices or systems without express written approval of NVIDIA Corporation.</p>
<p class="rubric-h2 rubric">Trademarks</p>
<p>NVIDIA and the NVIDIA logo are trademarks or registered trademarks of NVIDIA Corporation in the U.S. and other countries. Other company and product names may be trademarks of the respective companies with which they are associated.</p>
</section>
</section>
</div>
</div>
<footer>
<hr/>
<div role="contentinfo">
<p>© Copyright 2018-2024, NVIDIA Corporation & Affiliates. All rights reserved.
<span class="lastupdated">Last updated on Mar 06, 2024.
</span></p>
</div>
</footer>
</div>
</div>
</section>
</div>
<script>
jQuery(function () {
SphinxRtdTheme.Navigation.enable(true);
});
</script>
</body>
</html>
|
