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 2543 2544 2545 2546 2547 2548 2549 2550 2551 2552 2553 2554 2555 2556 2557 2558 2559 2560 2561 2562 2563 2564 2565 2566 2567 2568 2569 2570 2571 2572 2573 2574 2575 2576 2577 2578 2579 2580 2581 2582 2583 2584 2585 2586 2587 2588 2589 2590 2591 2592 2593 2594 2595 2596 2597 2598 2599 2600 2601 2602 2603 2604 2605 2606 2607 2608 2609 2610 2611 2612 2613 2614 2615 2616 2617 2618 2619 2620 2621 2622 2623 2624 2625 2626 2627 2628 2629 2630 2631 2632 2633 2634 2635 2636 2637 2638 2639 2640 2641 2642 2643 2644 2645 2646 2647 2648 2649 2650 2651 2652 2653 2654 2655 2656 2657 2658 2659 2660 2661 2662 2663 2664 2665 2666 2667 2668 2669 2670 2671 2672 2673 2674 2675 2676 2677 2678 2679 2680 2681 2682 2683 2684 2685 2686 2687 2688 2689 2690 2691 2692 2693 2694 2695 2696 2697 2698 2699 2700 2701 2702 2703 2704 2705 2706 2707 2708 2709 2710 2711 2712 2713 2714 2715 2716 2717 2718 2719 2720 2721 2722 2723 2724 2725 2726 2727 2728 2729 2730 2731 2732 2733 2734 2735 2736 2737 2738 2739 2740 2741 2742 2743 2744 2745 2746 2747 2748 2749 2750 2751 2752 2753 2754 2755 2756 2757 2758 2759 2760 2761 2762 2763 2764 2765 2766 2767 2768 2769 2770 2771 2772 2773 2774 2775 2776 2777 2778 2779 2780 2781 2782 2783 2784 2785 2786 2787 2788 2789 2790 2791 2792 2793 2794 2795 2796 2797 2798 2799 2800 2801 2802 2803 2804 2805 2806 2807 2808 2809 2810 2811 2812 2813 2814 2815 2816 2817 2818 2819 2820 2821 2822 2823 2824 2825 2826 2827 2828 2829 2830 2831 2832 2833 2834 2835 2836 2837 2838 2839 2840 2841 2842 2843 2844 2845 2846 2847 2848 2849 2850 2851 2852 2853 2854 2855 2856 2857 2858 2859 2860 2861 2862 2863 2864 2865 2866 2867 2868 2869 2870 2871 2872 2873 2874 2875 2876 2877 2878 2879 2880 2881 2882 2883 2884 2885 2886 2887 2888 2889 2890 2891 2892 2893 2894 2895 2896 2897 2898 2899 2900 2901 2902 2903 2904 2905 2906 2907 2908 2909 2910 2911 2912 2913 2914 2915 2916 2917 2918 2919 2920 2921 2922 2923 2924 2925 2926 2927 2928 2929 2930 2931 2932 2933 2934 2935 2936 2937 2938 2939 2940 2941 2942 2943 2944 2945 2946 2947 2948 2949 2950 2951 2952 2953 2954 2955 2956 2957 2958 2959 2960 2961 2962 2963 2964 2965 2966 2967 2968 2969 2970 2971 2972 2973 2974 2975 2976 2977 2978 2979 2980 2981 2982 2983 2984 2985 2986 2987 2988 2989 2990 2991 2992 2993 2994 2995 2996 2997 2998 2999 3000 3001 3002 3003 3004 3005 3006 3007 3008 3009 3010 3011 3012 3013 3014 3015 3016 3017 3018 3019 3020 3021 3022 3023 3024 3025 3026 3027 3028 3029 3030 3031 3032 3033 3034 3035 3036 3037 3038 3039 3040 3041 3042 3043 3044 3045 3046 3047 3048 3049 3050 3051 3052 3053 3054 3055 3056 3057 3058 3059 3060 3061 3062 3063 3064 3065 3066 3067 3068 3069 3070 3071 3072 3073 3074 3075 3076 3077 3078 3079 3080 3081 3082 3083 3084 3085 3086 3087 3088 3089 3090 3091 3092 3093 3094 3095 3096 3097 3098 3099 3100 3101 3102 3103 3104 3105 3106 3107 3108 3109 3110 3111 3112 3113 3114 3115 3116 3117 3118 3119 3120 3121 3122 3123 3124 3125 3126 3127 3128 3129 3130 3131 3132 3133 3134 3135 3136 3137 3138 3139 3140 3141 3142 3143 3144 3145 3146 3147 3148 3149 3150 3151 3152 3153 3154 3155 3156 3157 3158 3159 3160 3161 3162 3163 3164 3165 3166 3167 3168 3169 3170 3171 3172 3173 3174 3175 3176 3177 3178 3179 3180 3181 3182 3183 3184 3185 3186 3187 3188 3189 3190 3191 3192 3193 3194 3195 3196 3197 3198 3199 3200 3201 3202 3203 3204 3205 3206 3207 3208 3209 3210 3211 3212 3213 3214 3215 3216 3217 3218 3219 3220 3221 3222 3223 3224 3225 3226 3227 3228 3229 3230 3231 3232 3233 3234 3235 3236 3237 3238 3239 3240 3241 3242 3243 3244 3245 3246 3247 3248 3249 3250 3251 3252 3253 3254 3255 3256 3257 3258 3259 3260 3261 3262 3263 3264 3265 3266 3267 3268 3269 3270 3271 3272 3273 3274 3275 3276 3277 3278 3279 3280 3281 3282 3283 3284 3285 3286 3287 3288 3289 3290 3291 3292 3293 3294 3295 3296 3297 3298 3299 3300 3301 3302 3303 3304 3305 3306 3307 3308 3309 3310 3311 3312 3313 3314 3315 3316 3317 3318 3319 3320 3321 3322 3323 3324 3325 3326 3327 3328 3329 3330 3331 3332 3333 3334 3335 3336 3337 3338 3339 3340 3341 3342 3343 3344 3345 3346 3347 3348 3349 3350 3351 3352 3353 3354 3355 3356 3357 3358 3359 3360 3361 3362 3363 3364 3365 3366 3367 3368 3369 3370 3371 3372 3373 3374 3375 3376 3377 3378 3379 3380 3381 3382 3383 3384 3385 3386 3387 3388 3389 3390 3391 3392 3393 3394 3395 3396 3397 3398 3399 3400 3401 3402 3403 3404 3405 3406 3407 3408 3409 3410 3411 3412 3413 3414 3415 3416 3417 3418 3419 3420 3421 3422 3423 3424 3425 3426 3427 3428 3429 3430 3431 3432 3433 3434 3435 3436 3437 3438 3439 3440 3441 3442 3443 3444 3445 3446 3447 3448 3449 3450 3451 3452 3453 3454 3455 3456 3457 3458 3459 3460 3461 3462 3463 3464 3465 3466 3467 3468 3469 3470 3471 3472 3473 3474 3475 3476 3477 3478 3479 3480 3481 3482 3483 3484 3485 3486 3487 3488 3489 3490 3491 3492 3493 3494 3495 3496 3497 3498 3499 3500 3501 3502 3503 3504 3505 3506 3507 3508 3509 3510 3511 3512 3513 3514 3515 3516 3517 3518 3519 3520 3521 3522 3523 3524 3525 3526 3527 3528 3529 3530 3531 3532 3533 3534 3535 3536 3537 3538 3539 3540 3541 3542 3543 3544 3545 3546 3547 3548 3549 3550 3551 3552 3553 3554 3555 3556 3557 3558 3559 3560 3561 3562 3563 3564 3565 3566 3567 3568 3569 3570 3571 3572 3573 3574 3575 3576 3577 3578 3579 3580 3581 3582 3583 3584 3585 3586 3587 3588 3589 3590 3591 3592 3593 3594 3595 3596 3597 3598 3599 3600 3601 3602 3603 3604 3605 3606 3607 3608 3609 3610 3611 3612 3613 3614 3615 3616 3617 3618 3619 3620 3621 3622 3623 3624 3625 3626 3627 3628 3629 3630 3631 3632 3633 3634 3635 3636 3637 3638 3639 3640 3641 3642 3643 3644 3645 3646 3647 3648 3649 3650 3651 3652 3653 3654 3655 3656 3657 3658 3659 3660 3661 3662 3663 3664 3665 3666 3667 3668 3669 3670 3671 3672 3673 3674 3675 3676 3677 3678 3679 3680 3681 3682 3683 3684 3685 3686 3687 3688 3689 3690 3691 3692 3693 3694 3695 3696 3697 3698 3699 3700 3701 3702 3703 3704 3705 3706 3707 3708 3709 3710 3711 3712 3713 3714 3715 3716 3717 3718 3719 3720 3721 3722 3723 3724 3725 3726 3727 3728 3729 3730 3731 3732 3733 3734 3735 3736 3737 3738 3739 3740 3741 3742 3743 3744 3745 3746 3747 3748 3749 3750 3751 3752 3753 3754 3755 3756 3757 3758 3759 3760 3761 3762 3763 3764 3765 3766 3767 3768 3769 3770 3771 3772 3773 3774 3775 3776 3777 3778 3779 3780 3781 3782 3783 3784 3785 3786 3787 3788 3789 3790 3791 3792 3793 3794 3795 3796 3797 3798 3799 3800 3801 3802 3803 3804 3805 3806 3807 3808 3809 3810 3811 3812 3813 3814 3815 3816 3817 3818 3819 3820 3821 3822 3823 3824 3825 3826 3827 3828 3829 3830 3831 3832 3833 3834 3835 3836 3837 3838 3839 3840 3841 3842 3843 3844 3845 3846 3847 3848 3849 3850 3851 3852 3853 3854 3855 3856 3857 3858 3859 3860 3861 3862 3863 3864 3865 3866 3867 3868 3869 3870 3871 3872 3873 3874 3875 3876 3877 3878 3879 3880 3881 3882 3883 3884 3885 3886 3887 3888 3889 3890 3891 3892 3893 3894 3895 3896 3897 3898 3899 3900 3901 3902 3903 3904 3905 3906 3907 3908 3909 3910 3911 3912 3913 3914 3915 3916 3917 3918 3919 3920 3921 3922 3923 3924 3925 3926 3927 3928 3929 3930 3931 3932 3933 3934 3935 3936 3937 3938 3939 3940 3941 3942 3943 3944 3945 3946 3947 3948 3949 3950 3951 3952 3953 3954 3955 3956 3957 3958 3959 3960 3961 3962 3963 3964 3965 3966 3967 3968 3969 3970 3971 3972 3973 3974 3975 3976 3977 3978 3979 3980 3981 3982 3983
|
<!DOCTYPE html
PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" lang="en-us" xml:lang="en-us">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8"></meta>
<meta http-equiv="X-UA-Compatible" content="IE=edge"></meta>
<meta name="copyright" content="(C) Copyright 2005"></meta>
<meta name="DC.rights.owner" content="(C) Copyright 2005"></meta>
<meta name="DC.Type" content="concept"></meta>
<meta name="DC.Title" content="Kernel Profiling Guide"></meta>
<meta name="abstract" content="Nsight Compute profiling guide."></meta>
<meta name="description" content="Nsight Compute profiling guide."></meta>
<meta name="DC.Coverage" content="Nsight Compute"></meta>
<meta name="DC.subject" content="Kernel Profiling Guide"></meta>
<meta name="keywords" content="Kernel Profiling Guide"></meta>
<meta name="DC.Format" content="XHTML"></meta>
<meta name="DC.Identifier" content="abstract"></meta>
<link rel="stylesheet" type="text/css" href="../common/formatting/commonltr.css"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/site.css"></link>
<title>Kernel Profiling Guide :: Nsight Compute Documentation</title>
<!--[if lt IE 9]>
<script src="../common/formatting/html5shiv-printshiv.min.js"></script>
<![endif]-->
<script type="text/javascript" charset="utf-8" src="../common/scripts/tynt/tynt.js"></script>
-->
<script src="https://assets.adobedtm.com/5d4962a43b79/c1061d2c5e7b/launch-191c2462b890.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.ba-hashchange.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.scrollintoview.min.js"></script>
<script type="text/javascript" src="../search/htmlFileList.js"></script>
<script type="text/javascript" src="../search/htmlFileInfoList.js"></script>
<script type="text/javascript" src="../search/nwSearchFnt.min.js"></script>
<script type="text/javascript" src="../search/stemmers/en_stemmer.min.js"></script>
<script type="text/javascript" src="../search/index-1.js"></script>
<script type="text/javascript" src="../search/index-2.js"></script>
<script type="text/javascript" src="../search/index-3.js"></script>
<link rel="canonical" href="https://developer.nvidia.com/ProfilingGuide/index.html"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/qwcode.highlight.css"></link>
</head>
<body>
<header id="header"><span id="company">NVIDIA</span><span id="site-title">Nsight Compute Documentation</span><form id="search" method="get" action="search">
<input type="text" name="search-text"></input><fieldset id="search-location">
<legend>Search In:</legend>
<label><input type="radio" name="search-type" value="site"></input>Entire Site</label>
<label><input type="radio" name="search-type" value="document"></input>Just This Document</label></fieldset>
<button type="reset">clear search</button>
<button id="submit" type="submit">search</button></form>
</header>
<div id="site-content">
<nav id="site-nav">
<div class="category closed"><a href="../index.html" title="The root of the site.">Nsight Compute
v2022.3.0</a></div>
<div class="category"><a href="index.html" title="Kernel Profiling Guide">Kernel Profiling Guide</a></div>
<ul>
<li>
<div class="section-link"><a href="#introduction">1. Introduction</a></div>
<ul>
<li>
<div class="section-link"><a href="#profiling-applications">1.1. Profiling Applications</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#metric-collection">2. Metric Collection</a></div>
<ul>
<li>
<div class="section-link"><a href="#sets-and-sections">2.1. Sets and Sections</a></div>
</li>
<li>
<div class="section-link"><a href="#sections-and-rules">2.2. Sections and Rules</a></div>
</li>
<li>
<div class="section-link"><a href="#replay">2.3. Replay</a></div>
<ul>
<li>
<div class="section-link"><a href="#kernel-replay">2.3.1. Kernel Replay</a></div>
</li>
<li>
<div class="section-link"><a href="#application-replay">2.3.2. Application Replay</a></div>
</li>
<li>
<div class="section-link"><a href="#range-replay">2.3.3. Range Replay</a></div>
<ul>
<li>
<div class="section-link"><a href="#range-replay-define-range">2.3.3.1. Defining Ranges</a></div>
</li>
<li>
<div class="section-link"><a href="#range-replay-supported-apis">2.3.3.2. Supported APIs</a></div>
</li>
</ul>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#profile-series">2.4. Profile Series</a></div>
</li>
<li>
<div class="section-link"><a href="#overhead">2.5. Overhead</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#metrics-guide">3. Metrics Guide</a></div>
<ul>
<li>
<div class="section-link"><a href="#metrics-hw-model">3.1. Hardware Model</a></div>
</li>
<li>
<div class="section-link"><a href="#metrics-structure">3.2. Metrics Structure</a></div>
</li>
<li>
<div class="section-link"><a href="#metrics-decoder">3.3. Metrics Decoder</a></div>
</li>
<li>
<div class="section-link"><a href="#range-and-precision">3.4. Range and Precision</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#metrics-reference">4. Metrics Reference</a></div>
</li>
<li>
<div class="section-link"><a href="#sampling">5. Sampling</a></div>
<ul>
<li>
<div class="section-link"><a href="#statistical-sampler">5.1. Warp Scheduler States</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#reproducibility">6. Reproducibility</a></div>
<ul>
<li>
<div class="section-link"><a href="#serialization">6.1. Serialization</a></div>
</li>
<li>
<div class="section-link"><a href="#clock-control">6.2. Clock Control</a></div>
</li>
<li>
<div class="section-link"><a href="#cache-control">6.3. Cache Control</a></div>
</li>
<li>
<div class="section-link"><a href="#persistence-mode">6.4. Persistence Mode</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#special-configurations">7. Special Configurations</a></div>
<ul>
<li>
<div class="section-link"><a href="#special-configurations-mig">7.1. Multi Instance GPU</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#roofline">8. Roofline Charts</a></div>
<ul>
<li>
<div class="section-link"><a href="#roofline-overview">8.1. Overview</a></div>
</li>
<li>
<div class="section-link"><a href="#roofline-analysis">8.2. Analysis</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#memory-chart">9. Memory Chart</a></div>
<ul>
<li>
<div class="section-link"><a href="#memory-chart-overview">9.1. Overview</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#memory-tables">10. Memory Tables</a></div>
<ul>
<li>
<div class="section-link"><a href="#memory-tables-smem">10.1. Shared Memory</a></div>
</li>
<li>
<div class="section-link"><a href="#memory-tables-l1">10.2. L1/TEX Cache</a></div>
</li>
<li>
<div class="section-link"><a href="#memory-tables-l2">10.3. L2 Cache</a></div>
</li>
<li>
<div class="section-link"><a href="#memory-tables-l2-evict-policy">10.4. L2 Cache Eviction Policies</a></div>
</li>
<li>
<div class="section-link"><a href="#memory-tables-dram">10.5. Device Memory</a></div>
</li>
</ul>
</li>
<li>
<div class="section-link"><a href="#faq">11. FAQ</a></div>
</li>
</ul>
</nav>
<div id="resize-nav"></div>
<nav id="search-results">
<h2>Search Results</h2>
<ol></ol>
</nav>
<div id="contents-container">
<div id="breadcrumbs-container">
<div id="release-info">Kernel Profiling Guide
(<a href="../pdf/ProfilingGuide.pdf">PDF</a>)
-
v2022.3.0
(<a href="https://developer.nvidia.com">older</a>)
-
Last updated August 24, 2022
-
<a href="mailto:devtools@nvidia.com?subject=Nsight Compute Documentation Feedback: Kernel Profiling Guide">Send Feedback</a>
-
<span class="st_facebook"></span><span class="st_twitter"></span><span class="st_linkedin"></span><span class="st_reddit"></span><span class="st_slashdot"></span><span class="st_tumblr"></span><span class="st_sharethis"></span></div>
</div>
<article id="contents">
<div class="topic nested0" id="abstract"><a name="abstract" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#abstract" name="abstract" shape="rect">Kernel Profiling Guide</a></h2>
<div class="body conbody">
<p class="shortdesc">Nsight Compute profiling guide.</p>
<p class="p">
Kernel Profiling Guide with metric types and meaning, data collection modes and FAQ for common problems.
</p>
</div>
</div>
<div class="topic concept nested0" id="introduction"><a name="introduction" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#introduction" name="introduction" shape="rect">1. Introduction</a></h2>
<div class="body conbody">
<p class="p">
This guide describes various profiling topics related to <span class="keyword">NVIDIA Nsight Compute</span> and <span class="keyword">NVIDIA Nsight Compute CLI</span>.
Most of these apply to both the UI and the CLI version of the tool.
</p>
<div class="p">
To use the tools effectively, it is recommended to read this guide, as well as at least the following chapters of the <dfn class="term">CUDA Programming Guide</dfn>:
<ul class="ul">
<li class="li"><a class="xref" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programming-model" target="_blank" shape="rect">Programming Model</a></li>
<li class="li"><a class="xref" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#hardware-implementation" target="_blank" shape="rect">Hardware Implementation</a></li>
<li class="li"><a class="xref" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#performance-guidelines" target="_blank" shape="rect">Performance Guidelines</a></li>
</ul>
Afterwards, it should be enough to read the <dfn class="term">Quickstart</dfn> chapter of the <span class="keyword">NVIDIA Nsight Compute</span> or <span class="keyword">NVIDIA Nsight Compute CLI</span>
documentation, respectively, to start using the tools.
</div>
</div>
<div class="topic concept nested1" id="profiling-applications"><a name="profiling-applications" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#profiling-applications" name="profiling-applications" shape="rect">1.1. Profiling Applications</a></h3>
<div class="body conbody">
<p class="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>
<div class="fig fignone"><span class="desc figdesc">Regular Application Execution</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/regular-application-process.png"></img></div><br clear="none"></br></div>
<p class="p">
When profiling an application with <span class="keyword">NVIDIA Nsight Compute</span>, the behavior is different.
The user launches the <span class="keyword">NVIDIA Nsight Compute</span> 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 class="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>
<div class="fig fignone"><span class="desc figdesc">Profiled Application Execution</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/profiled-process.png"></img></div><br clear="none"></br></div>
</div>
</div>
</div>
<div class="topic concept nested0" id="metric-collection"><a name="metric-collection" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#metric-collection" name="metric-collection" shape="rect">2. Metric Collection</a></h2>
<div class="body conbody">
<p class="p">
Collection of performance metrics is the key feature of <span class="keyword">NVIDIA Nsight Compute</span>.
Since there is a huge list of metrics available, it is often easier to use some of the tool's
pre-defined <a class="xref" href="index.html#sets-and-sections" shape="rect">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="xref" href="index.html#overhead" shape="rect">Overhead</a> associated with data collection.
</p>
</div>
<div class="topic concept nested1" id="sets-and-sections"><a name="sets-and-sections" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#sets-and-sections" name="sets-and-sections" shape="rect">2.1. Sets and Sections</a></h3>
<div class="body conbody">
<p class="p"><span class="keyword">NVIDIA Nsight Compute</span> uses <dfn class="term">Section Sets</dfn> (short <dfn class="term">sets</dfn>) to decide,
on a very high level, the amount of metrics to be collected.
Each set includes one or more <dfn class="term">Sections</dfn>, 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 class="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="xref" href="index.html#overhead" shape="rect">Overhead</a> for more information on profiling overhead.
</p>
<p class="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 default set is collected when no <samp class="ph codeph">--set</samp>, <samp class="ph codeph">--section</samp> and no <samp class="ph codeph">--metrics</samp>
options are passed on the command line. The full set of sections can be collected with <samp class="ph codeph">--set full</samp>.
</p>
<p class="p">
Use <samp class="ph codeph">--list-sets</samp> to see the list of currently available sets.
Use <samp class="ph codeph">--list-sections</samp> to see the list of currently available sections.
The default search directory and the location of pre-defined section files are also called <samp class="ph codeph">sections/</samp>.
All related command line options can be found in the <span class="keyword">NVIDIA Nsight Compute CLI</span> documentation.
</p>
</div>
</div>
<div class="topic concept nested1" id="sections-and-rules"><a name="sections-and-rules" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#sections-and-rules" name="sections-and-rules" shape="rect">2.2. Sections and Rules</a></h3>
<div class="body conbody">
<div class="p">
<div class="tablenoborder"><a name="sections-and-rules__available-sections" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="sections-and-rules__available-sections" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 1. Available Sections</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e239" rowspan="1" colspan="1">Identifier and Filename</th>
<th class="entry" valign="top" width="66.66666666666666%" id="d54e242" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">ComputeWorkloadAnalysis (Compute Workload Analysis)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">InstructionStats (Instruction Statistics)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">LaunchStats (Launch Statistics)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">MemoryWorkloadAnalysis (Memory Workload Analysis)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">Nvlink (Nvlink)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
High-level summary of NVLink utilization. It shows the total received and transmitted (sent) memory, as well as the overall
link peak utilization.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">Nvlink_Tables (Nvlink_Tables)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
Detailed tables with properties for each NVLink.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">Nvlink_Topology (Nvlink_Topology)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
NVLink Topology diagram shows logical NVLink connections with transmit/receive throughput.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">Occupancy (Occupancy)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">SchedulerStats (Scheduler Statistics)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">SourceCounters (Source Counters)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">SpeedOfLight (GPU Speed Of Light Throughput)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e239" rowspan="1" colspan="1">WarpStateStats (Warp State Statistics)</td>
<td class="entry" valign="top" width="66.66666666666666%" headers="d54e242" rowspan="1" colspan="1">
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.
</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested1" id="replay"><a name="replay" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#replay" name="replay" shape="rect">2.3. Replay</a></h3>
<div class="body conbody">
<p class="p">
Depending on which metrics are to be collected, kernels might need to be <dfn class="term">replayed</dfn>
one or more times, since not all metrics can be collected in a single <dfn class="term">pass</dfn>.
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>
</div>
<div class="topic concept nested2" id="kernel-replay"><a name="kernel-replay" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#kernel-replay" name="kernel-replay" shape="rect">2.3.1. Kernel Replay</a></h3>
<div class="body conbody">
<p class="p">
In <dfn class="term">Kernel Replay</dfn>, all metrics requested for a specific kernel instance in <span class="keyword">NVIDIA Nsight Compute</span> 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 class="p"><span class="keyword">NVIDIA Nsight Compute</span> 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 class="p">
As explained in <a class="xref" href="index.html#overhead" shape="rect">Overhead</a>, the time needed for this increases the more memory is accessed, especially written, by a kernel.
If <span class="keyword">NVIDIA Nsight Compute</span> 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>
<div class="fig fignone"><span class="desc figdesc">Regular Application Execution</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/replay-regular-execution.png"></img></div><br clear="none"></br></div>
<div class="fig fignone"><span class="desc figdesc">Execution with Kernel Replay. All memory is saved, and memory written by the kernel is restored in-between replay passes.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/replay-kernel.png"></img></div><br clear="none"></br></div>
</div>
</div>
<div class="topic concept nested2" id="application-replay"><a name="application-replay" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#application-replay" name="application-replay" shape="rect">2.3.2. Application Replay</a></h3>
<div class="body conbody">
<p class="p">
In <dfn class="term">Application Replay</dfn>, all metrics requested for a specific kernel launch in <span class="keyword">NVIDIA Nsight Compute</span> are grouped into one or more passes.
In contrast to <a class="xref" href="index.html#kernel-replay" shape="rect">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 class="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 class="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="xref" href="index.html#cache-control" shape="rect">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 class="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 class="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>
<div class="fig fignone"><span class="desc figdesc">Regular Application Execution</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/replay-regular-execution.png"></img></div><br clear="none"></br></div>
<div class="fig fignone"><span class="desc figdesc">Execution with Application Replay. No memory is saved or restored, but the cost of running the application itself is duplicated.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/replay-application.png"></img></div><br clear="none"></br></div>
<p class="p">
Across application replay passes, <span class="keyword">NVIDIA Nsight Compute</span> matches metric data for the individual, selected kernel launches.
The matching strategy can be selected using the <samp class="ph codeph">--app-replay-match</samp> option.
For matching, only kernels within the same process and running on the same device are considered.
By default, the <dfn class="term">grid</dfn> 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>
<div class="fig fignone"><span class="desc figdesc">Kernel matching during application replay using the <dfn class="term">grid</dfn> strategy.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/replay-application-kernel-matching.png"></img></div><br clear="none"></br></div>
</div>
</div>
<div class="topic concept nested2" id="range-replay"><a name="range-replay" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#range-replay" name="range-replay" shape="rect">2.3.3. Range Replay</a></h3>
<div class="body conbody">
<p class="p">
In <dfn class="term">Range Replay</dfn>, all requested metrics in <span class="keyword">NVIDIA Nsight Compute</span> are grouped into one or more passes.
In contrast to <a class="xref" href="index.html#kernel-replay" shape="rect">Kernel Replay</a> and <a class="xref" href="index.html#application-replay" shape="rect">Application Replay</a>, <dfn class="term">Range Replay</dfn>
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>
<div class="fig fignone"><span class="desc figdesc">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><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/replay-range.png"></img></div><br clear="none"></br></div>
</div>
<div class="topic concept nested3" id="range-replay-define-range"><a name="range-replay-define-range" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#range-replay-define-range" name="range-replay-define-range" shape="rect">2.3.3.1. Defining Ranges</a></h3>
<div class="body conbody">
<p class="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 class="ul">
<li class="li"><strong class="ph b">Profiler Start/Stop API</strong><p class="p">
Set the start marker using <samp class="ph codeph">cu(da)ProfilerStart</samp> and the end marker using <samp class="ph codeph">cu(da)ProfilerStop</samp>.
Note: The CUDA driver API variants of this API require to include <samp class="ph codeph">cudaProfiler.h</samp>.
The CUDA runtime variants require to include <samp class="ph codeph">cuda_profiler_api.h</samp>.
</p>
<p class="p">
This is the default for <span class="keyword">NVIDIA Nsight Compute</span>.
</p>
</li>
<li class="li"><strong class="ph b">NVTX Ranges</strong><p class="p">
Define the range using an <a class="xref" href="../NsightComputeCli/index.html#nvtx-filtering" shape="rect">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 class="p">
The application must have been instrumented with the NVTX API for any expressions to match.
</p>
<p class="p">
This mode is enabled by passing <samp class="ph codeph">--nvtx --nvtx-include <expression> [--nvtx-include <expression>]</samp>
to the <span class="keyword">NVIDIA Nsight Compute CLI</span>.
</p>
</li>
</ul>
<div class="p">
Ranges must fulfill several requirements:
<ul class="ul">
<li class="li">It must be possible to synchronize all active CUDA contexts at the start of the range.</li>
<li class="li">Ranges must not include unsupported CUDA API calls. See <a class="xref" href="index.html#range-replay-supported-apis" shape="rect">Supported APIs</a> for the list of currently supported APIs.
</li>
</ul>
In addition, there are several recommendations that ranges should comply with to guarantee a correct capture and replay:
<ul class="ul">
<li class="li">
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.
</li>
<li class="li">
Avoid freeing host allocations written by device memory during the range.
This includes both heap as well as stack allocations.
<span class="keyword">NVIDIA Nsight Compute</span> 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 <samp class="ph codeph">--launch-count 1</samp>,
set the <dfn class="term">disable-host-restore</dfn> range replay option and optionally use <samp class="ph codeph">--kill yes</samp> to terminate the process after this range.
</li>
<li class="li">
When defining the range markers using <samp class="ph codeph">cu(da)ProfilerStart/Stop</samp>, prefer the CUDA driver API calls <samp class="ph codeph">cuProfilerStart/Stop</samp>.
Internally, <span class="keyword">NVIDIA Nsight Compute</span> 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.
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested3" id="range-replay-supported-apis"><a name="range-replay-supported-apis" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#range-replay-supported-apis" name="range-replay-supported-apis" shape="rect">2.3.3.2. Supported APIs</a></h3>
<div class="body conbody">
<p class="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="xref" href="https://docs.nvidia.com/cuda/cuda-driver-api/index.html" target="_blank" shape="rect">CUDA Driver API documentation</a>.
</p>
<p class="p">
Generally, range replay only captures and replay CUDA <dfn class="term">Driver</dfn> API calls.
CUDA <dfn class="term">Runtime</dfn> APIs calls can be captured when they generate only supported CUDA Driver API calls internally.
Deprecated APIs are not supported.
</p>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-error-handling"><a name="range-replay-supported-apis__rr-supported-apis-error-handling" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Error Handling</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-initialization"><a name="range-replay-supported-apis__rr-supported-apis-initialization" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Initialization</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-version"><a name="range-replay-supported-apis__rr-supported-apis-version" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Version Management</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-device"><a name="range-replay-supported-apis__rr-supported-apis-device" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Device Management</h5>
<div class="p">
All supported, except:
<ul class="ul">
<li class="li">cuDeviceSetMemPool</li>
</ul>
</div>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-primary-ctx"><a name="range-replay-supported-apis__rr-supported-apis-primary-ctx" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Primary Context Management</h5>
<div class="p">
<ul class="ul">
<li class="li">cuDevicePrimaryCtxGetState</li>
</ul>
</div>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-ctx"><a name="range-replay-supported-apis__rr-supported-apis-ctx" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Context Management</h5>
<div class="p">
All supported, except:
<ul class="ul">
<li class="li">cuCtxSetCacheConfig</li>
<li class="li">cuCtxSetSharedMemConfig</li>
</ul>
</div>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-module"><a name="range-replay-supported-apis__rr-supported-apis-module" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Module Management</h5>
<div class="p">
<ul class="ul">
<li class="li">cuModuleGetFunction</li>
<li class="li">cuModuleGetGlobal</li>
<li class="li">cuModuleGetSurfRef</li>
<li class="li">cuModuleGetTexRef</li>
<li class="li">cuModuleLoad</li>
<li class="li">cuModuleLoadData</li>
<li class="li">cuModuleLoadDataEx</li>
<li class="li">cuModuleLoadFatBinary</li>
<li class="li">cuModuleUnload</li>
</ul>
</div>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-memory"><a name="range-replay-supported-apis__rr-supported-apis-memory" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Memory Management</h5>
<div class="p">
<ul class="ul">
<li class="li">cuArray*</li>
<li class="li">cuDeviceGetByPCIBusId</li>
<li class="li">cuDeviceGetPCIBusId</li>
<li class="li">cuMemAlloc</li>
<li class="li">cuMemAllocHost</li>
<li class="li">cuMemAllocPitch</li>
<li class="li">cuMemcpy*</li>
<li class="li">cuMemFree</li>
<li class="li">cuMemFreeHost</li>
<li class="li">cuMemGetAddressRange</li>
<li class="li">cuMemGetInfo</li>
<li class="li">cuMemHostAlloc</li>
<li class="li">cuMemHostGetDevicePointer</li>
<li class="li">cuMemHostGetFlags</li>
<li class="li">cuMemHostRegister</li>
<li class="li">cuMemHostUnregister</li>
<li class="li">cuMemset*</li>
<li class="li">cuMipmapped*</li>
</ul>
</div>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-virtual"><a name="range-replay-supported-apis__rr-supported-apis-virtual" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Virtual Memory Management</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-stream-ordered"><a name="range-replay-supported-apis__rr-supported-apis-stream-ordered" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Stream Ordered Memory Allocator</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-unified-addr"><a name="range-replay-supported-apis__rr-supported-apis-unified-addr" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Unified Addressing</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-stream"><a name="range-replay-supported-apis__rr-supported-apis-stream" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Stream Management</h5>
<div class="p">
<ul class="ul">
<li class="li">cuStreamCreate*</li>
<li class="li">cuStreamDestroy</li>
<li class="li">cuStreamGet*</li>
<li class="li">cuStreamQuery</li>
<li class="li">cuStreamSetAttribute</li>
<li class="li">cuStreamSynchronize</li>
<li class="li">cuStreamWaitEvent</li>
</ul>
</div>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-event"><a name="range-replay-supported-apis__rr-supported-apis-event" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Event Management</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-externa-resource"><a name="range-replay-supported-apis__rr-supported-apis-externa-resource" shape="rect">
<!-- --></a><h5 class="title sectiontitle">External Resource interoperability</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-stream-memory"><a name="range-replay-supported-apis__rr-supported-apis-stream-memory" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Stream Memory Operations</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-execution"><a name="range-replay-supported-apis__rr-supported-apis-execution" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Execution Control</h5>
<div class="p">
<ul class="ul">
<li class="li">cuFuncGetAttribute</li>
<li class="li">cuFuncGetModule</li>
<li class="li">cuFuncSetAttribute</li>
<li class="li">cuFuncSetCacheConfig</li>
<li class="li">cuLaunchCooperativeKernel</li>
<li class="li">cuLaunchHostFunc</li>
<li class="li">cuLaunchKernel</li>
</ul>
</div>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-graph"><a name="range-replay-supported-apis__rr-supported-apis-graph" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Graph Management</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-occupancy"><a name="range-replay-supported-apis__rr-supported-apis-occupancy" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Occupancy</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-tex-surf"><a name="range-replay-supported-apis__rr-supported-apis-tex-surf" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Texture/Surface Reference Management</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-tex"><a name="range-replay-supported-apis__rr-supported-apis-tex" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Texture Object Management</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-surf"><a name="range-replay-supported-apis__rr-supported-apis-surf" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Surface Object Management</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-peer"><a name="range-replay-supported-apis__rr-supported-apis-peer" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Peer Context Memory Access</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-graphics"><a name="range-replay-supported-apis__rr-supported-apis-graphics" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Graphics Interoperability</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-driver-entry"><a name="range-replay-supported-apis__rr-supported-apis-driver-entry" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Driver Entry Point Access</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-surf"><a name="range-replay-supported-apis__rr-supported-apis-surf" shape="rect">
<!-- --></a><h5 class="title sectiontitle">Surface Object Management</h5>
<p class="p">
All supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-opengl"><a name="range-replay-supported-apis__rr-supported-apis-opengl" shape="rect">
<!-- --></a><h5 class="title sectiontitle">OpenGL Interoperability</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-vdpau"><a name="range-replay-supported-apis__rr-supported-apis-vdpau" shape="rect">
<!-- --></a><h5 class="title sectiontitle">VDPAU Interoperability</h5>
<p class="p">
Not supported.
</p>
</div>
<div class="section" id="range-replay-supported-apis__rr-supported-apis-egl"><a name="range-replay-supported-apis__rr-supported-apis-egl" shape="rect">
<!-- --></a><h5 class="title sectiontitle">EGL Interoperability</h5>
<p class="p">
Not supported.
</p>
</div>
</div>
</div>
</div>
</div>
<div class="topic concept nested1" id="profile-series"><a name="profile-series" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#profile-series" name="profile-series" shape="rect">2.4. Profile Series</a></h3>
<div class="body conbody">
<p class="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 class="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>
<div class="fig fignone"><span class="desc figdesc">Profile Series action.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/profile-series-action.png"></img></div><br clear="none"></br></div>
<div class="fig fignone"><span class="desc figdesc">Profile Series dialog.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/profile-series-dialog.png"></img></div><br clear="none"></br></div>
</div>
</div>
<div class="topic concept nested1" id="overhead"><a name="overhead" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#overhead" name="overhead" shape="rect">2.5. Overhead</a></h3>
<div class="body conbody">
<p class="p">
As with most measurements, collecting performance data using <span class="keyword">NVIDIA Nsight Compute CLI</span> incurs some runtime overhead on the application.
The overhead does depend on a number of different factors:
</p>
<ul class="ul">
<li class="li"><strong class="ph b">Number and type of collected metrics</strong><p class="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 class="p">
Furthermore, only a limited number of metrics can be collected in a single <dfn class="term">pass</dfn>
of the kernel execution. If more metrics are requested, the kernel launch is <dfn class="term">replayed</dfn>
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 class="li"><strong class="ph b">The collected section set</strong><p class="p">
Since each <a class="xref" href="index.html#sets-and-sections" shape="rect">set</a> specifies a group of section to be collected,
choosing a less comprehensive set can reduce profiling overhead.
See the <samp class="ph codeph">--set</samp> command in the
<a class="xref" href="../NsightComputeCli/index.html#command-line-options-profile" shape="rect"><span class="keyword">NVIDIA Nsight Compute CLI</span></a> documentation.
</p>
</li>
<li class="li"><strong class="ph b">Number of collected sections</strong><p class="p">
Since each <a class="xref" href="index.html#sets-and-sections" shape="rect">section</a> specifies a set metrics to be collected,
selecting fewer sections can reduce profiling overhead.
See the <samp class="ph codeph">--section</samp> command in the
<a class="xref" href="../NsightComputeCli/index.html#command-line-options-profile" shape="rect"><span class="keyword">NVIDIA Nsight Compute CLI</span></a> documentation.
</p>
</li>
<li class="li"><strong class="ph b">Number of profiled kernels</strong><p class="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="xref" href="../NsightComputeCli/index.html#command-line-options-profile" shape="rect"><span class="keyword">NVIDIA Nsight Compute CLI</span></a> documentation.
</p>
<p class="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 class="li"><strong class="ph b">GPU Architecture</strong><p class="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>
</div>
</div>
</div>
<div class="topic concept nested0" id="metrics-guide"><a name="metrics-guide" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#metrics-guide" name="metrics-guide" shape="rect">3. Metrics Guide</a></h2>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested1" id="metrics-hw-model"><a name="metrics-hw-model" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#metrics-hw-model" name="metrics-hw-model" shape="rect">3.1. Hardware Model</a></h3>
<div class="body conbody">
<div class="section" id="metrics-hw-model__metrics-hw-compute-model"><a name="metrics-hw-model__metrics-hw-compute-model" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Compute Model</h3>
<p class="p">
All NVIDIA GPUs are designed to support a general purpose heterogeneous
parallel programming model, commonly known as <dfn class="term">Compute</dfn>.
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>
<div class="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:
<ul class="ul">
<li class="li">A <dfn class="term">Grid</dfn> is a 1D, 2D or 3D array of thread blocks.
</li>
<li class="li">A <dfn class="term">Block</dfn> is a 1D, 2D or 3D array of threads, also known as a <dfn class="term">Cooperative Thread Array (CTA)</dfn>.
</li>
<li class="li">A <dfn class="term">Thread</dfn> is a single thread which runs on one of the GPU's SM units.
</li>
</ul>
</div>
<p class="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="xref" href="index.html#sections-and-rules" shape="rect">Launch Statistics</a> section.
</p>
<p class="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 <dfn class="term">occupancy</dfn>, and these physical resources limit
this occupancy.
Details on the kernel's occupancy are collected by the
<a class="xref" href="index.html#sections-and-rules" shape="rect">Occupancy</a> section.
</p>
<p class="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 class="p">
CTAs are further divided into groups of 32 threads called <dfn class="term">Warps</dfn>.
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 class="p">
The total number of CTAs that can run concurrently on a given GPU is referred to as <dfn class="term">Wave</dfn>.
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>
</div>
<div class="section" id="metrics-hw-model__metrics-hw-sm"><a name="metrics-hw-model__metrics-hw-sm" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Streaming Multiprocessor</h3>
<p class="p">
The <dfn class="term">Streaming Multiprocessor (SM)</dfn> 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 class="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="xref" href="index.html#sections-and-rules" shape="rect">Source Counters</a>
section allows you to inspect instruction execution and predication
details on the <dfn class="term">Source Page</dfn>, along with
<a class="xref" href="index.html#sampling" shape="rect">Sampling</a> information.
</p>
<div class="p">
Each SM is partitioned into four processing blocks, called <dfn class="term">SM sub partitions</dfn>.
The SM sub partitions are the primary processing elements on the SM.
Each sub partition contains the following units:
<ul class="ul">
<li class="li">Warp Scheduler</li>
<li class="li">Register File</li>
<li class="li">
Execution Units/Pipelines/Cores
<ul class="ul">
<li class="li">Integer Execution units</li>
<li class="li">Floating Point Execution units</li>
<li class="li">Memory Load/Store units</li>
<li class="li">Special Function unit</li>
<li class="li">Tensor Cores</li>
</ul>
</li>
</ul>
Shared within an SM across the four SM partitions are:
<ul class="ul">
<li class="li">Unified L1 Data Cache / Shared Memory</li>
<li class="li">Texture units</li>
<li class="li">RT Cores, if available</li>
</ul>
</div>
<p class="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 <dfn class="term">active</dfn> or <dfn class="term">resident</dfn>
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 <dfn class="term">eligible</dfn> 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="xref" href="index.html#sections-and-rules" shape="rect">Scheduler Statistics</a> section.
</p>
<div class="p">
A warp is <dfn class="term">stalled</dfn> when the warp is waiting on
<ul class="ul">
<li class="li">an instruction fetch,</li>
<li class="li">a memory dependency (result of memory instruction),</li>
<li class="li">an execution dependency (result of previous instruction), or</li>
<li class="li">a synchronization barrier.</li>
</ul>
See <a class="xref" href="index.html#statistical-sampler" shape="rect">Warp Scheduler States</a> for the list of stall reasons that can be profiled
and the <a class="xref" href="index.html#sections-and-rules" shape="rect">Warp State Statistics</a> section
for a summary of warp states found in the kernel execution.
</div>
<p class="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="xref" href="index.html#sections-and-rules" shape="rect">Launch Statistics</a> section shows
the kernel's register usage.
</p>
</div>
<div class="section" id="metrics-hw-model__metrics-hw-memory"><a name="metrics-hw-model__metrics-hw-memory" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Memory</h3>
<p class="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 class="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 class="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 class="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 class="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 class="p">
Detailed memory metrics are collected by the
<a class="xref" href="index.html#sections-and-rules" shape="rect">Memory Workload Analysis</a> section.
</p>
</div>
<div class="section" id="metrics-hw-model__metrics-hw-caches"><a name="metrics-hw-model__metrics-hw-caches" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Caches</h3>
<p class="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>
<div class="p">
<div class="fig fignone"><span class="desc figdesc">Model of the L2 cache.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/hw-model-lts.png"></img></div><br clear="none"></br></div>
</div>
<p class="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="xref" href="index.html#metrics-hw-model__metrics-hw-tex-surf" shape="rect">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 class="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 class="p">
Cache hit and miss rates as well as data transfers are reported in the
<a class="xref" href="index.html#sections-and-rules" shape="rect">Memory Workload Analysis</a> section.
</p>
<div class="p">
<div class="fig fignone"><span class="desc figdesc">Model of Load/Store and Texture pipelines for the L1TEX cache.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/hw-model-l1tex.png"></img></div><br clear="none"></br></div>
</div>
</div>
<div class="section" id="metrics-hw-model__metrics-hw-tex-surf"><a name="metrics-hw-model__metrics-hw-tex-surf" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Texture/Surface</h3>
<p class="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 class="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 class="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 class="p">
Information on texture and surface memory can be found in the
<a class="xref" href="index.html#sections-and-rules" shape="rect">Memory Workload Analysis</a> section.
</p>
</div>
</div>
</div>
<div class="topic concept nested1" id="metrics-structure"><a name="metrics-structure" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#metrics-structure" name="metrics-structure" shape="rect">3.2. Metrics Structure</a></h3>
<div class="body conbody">
<div class="section" id="metrics-structure__metrics-overview"><a name="metrics-structure__metrics-overview" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Metrics Overview</h3>
<p class="p"><span class="keyword">NVIDIA Nsight Compute</span> 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 class="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 class="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>
</div>
<div class="section" id="metrics-structure__metrics-entities"><a name="metrics-structure__metrics-entities" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Metrics Entities</h3>
<p class="p">
While in <span class="keyword">NVIDIA Nsight Compute</span>, all performance counters are named <dfn class="term">metrics</dfn>, they can be split further
into groups with specific properties.
For metrics collected via the <dfn class="term">PerfWorks</dfn> measurement library, the following entities exist:
</p>
<div class="p"><strong class="ph b">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 <dfn class="term">roll-ups</dfn>:
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.sum</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">The sum of counter values across all unit instances. </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.avg</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">The average counter value across all unit instances. </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.min</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">The minimum counter value across all unit instances. </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.max</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">The maximum counter value across all unit instances. </td>
</tr>
</table>
</div>
<div class="p">
Counter roll-ups have the following calculated quantities as built-in sub-metrics:
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_active </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate during unit active cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_active.per_second </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate during unit active cycles, per second * </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_elapsed </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate during unit elapsed cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_elapsed.per_second </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate during unit elapsed cycles, per second * </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_region </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate over a user-specified "range" </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_region.per_second </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate over a user-specified "range", per second * </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_frame </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate over a user-specified "frame" </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_sustained_frame.per_second </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak sustained rate over a user-specified "frame", per second * </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.per_second </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the number of operations per second </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.per_cycle_active </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the number of operations per unit active cycle </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.per_cycle_elapsed </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the number of operations per unit elapsed cycle </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.per_cycle_in_region </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the number of operations per user-specified "range" cycle </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.per_cycle_in_frame </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the number of operations per user-specified "frame" cycle </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_active </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak sustained rate achieved during unit active cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_elapsed </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak sustained rate achieved during unit elapsed cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_region </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak sustained rate achieved over a user-specified "range" </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_frame </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak sustained rate achieved over a user-specified "frame" </td>
</tr>
</table>
* sub-metrics added in <span class="keyword">NVIDIA Nsight Compute</span> 2022.2.0.
</div>
<p class="p">
Example: <samp class="ph codeph"><span class="keyword">ncu</span> --query-metrics-mode suffix --metrics sm__inst_executed --chip ga100</samp></p>
<div class="p"><strong class="ph b">Ratios</strong> have three sub-metrics:
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">The value expressed as a percentage. </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.ratio</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">The value expressed as a ratio. </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.max_rate</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">The ratio's maximum value. </td>
</tr>
</table>
</div>
<p class="p">
Example: <samp class="ph codeph"><span class="keyword">ncu</span> --query-metrics-mode suffix --metrics smsp__average_warp_latency --chip ga100</samp></p>
<div class="p"><strong class="ph b">Throughputs</strong> indicate how close a portion of the GPU reached to peak rate. Every throughput has the following sub-metrics:
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_active</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak sustained rate achieved during unit active cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_elapsed</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak sustained rate achieved during unit elapsed cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_region</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak sustained rate achieved over a user-specified "range" time </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_sustained_frame</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak sustained rate achieved over a user-specified "frame" time </td>
</tr>
</table>
</div>
<p class="p">
Example: <samp class="ph codeph"><span class="keyword">ncu</span> --query-metrics-mode suffix --metrics sm__throughput --chip ga100</samp></p>
<p class="p">
Throughputs have a breakdown of underlying metrics from which the throughput value is computed.
You can collect <samp class="ph codeph">breakdown:<throughput-metric></samp> to collect a throughput's breakdown metrics.
</p>
<div class="p"><strong class="ph b">Deprecated counter sub-metrics:</strong>
The following sub-metrics were removed in <span class="keyword">NVIDIA Nsight Compute</span> 2022.2.0, due to not being useful for performance optimization:
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.peak_burst </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">the peak burst rate </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_active </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak burst rate achieved during unit active cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_elapsed </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak burst rate achieved during unit elapsed cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_region </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak burst rate achieved over a user-specified "range" </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_frame </tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">% of peak burst rate achieved over a user-specified "frame" </td>
</tr>
</table>
</div>
<div class="p"><strong class="ph b">Deprecated throughput sub-metrics:</strong>
The following sub-metrics were removed in <span class="keyword">NVIDIA Nsight Compute</span> 2022.2, due to not being useful for performance optimization:
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_active</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak burst rate achieved during unit active cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_elapsed</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak burst rate achieved during unit elapsed cycles </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_region</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak burst rate achieved over a user-specified "range" time </td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">.pct_of_peak_burst_frame</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"> % of peak burst rate achieved over a user-specified "frame" time </td>
</tr>
</table>
</div>
<p class="p">
In addition to PerfWorks metrics, <span class="keyword">NVIDIA Nsight Compute</span> uses several other measurement providers that each generate their own metrics.
These are explained in the <a class="xref" href="index.html#metrics-reference" shape="rect">Metrics Reference</a>.
</p>
</div>
<div class="section" id="metrics-structure__metrics-examples"><a name="metrics-structure__metrics-examples" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Metrics Examples</h3>
<div class="p"><pre xml:space="preserve">
## 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.peak_sustained_frame
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained_frame.per_second
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained_region
l1tex__data_bank_conflicts_pipe_lsu.sum.peak_sustained_region.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_cycle_in_frame
l1tex__data_bank_conflicts_pipe_lsu.sum.per_cycle_in_region
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
l1tex__data_bank_conflicts_pipe_lsu.sum.pct_of_peak_sustained_frame
l1tex__data_bank_conflicts_pipe_lsu.sum.pct_of_peak_sustained_region
...
</pre></div>
</div>
<div class="section" id="metrics-structure__metrics-naming"><a name="metrics-structure__metrics-naming" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Metrics Naming Conventions</h3>
<div class="p">
Counters and metrics _generally_ obey the naming scheme:
<ul class="ul">
<li class="li"> Unit-Level Counter : <tt class="ph tt">unit__(subunit?)_(pipestage?)_quantity_(qualifiers?)</tt></li>
<li class="li"> Interface Counter : <tt class="ph tt">unit__(subunit?)_(pipestage?)_(interface)_quantity_(qualifiers?)</tt></li>
<li class="li"> Unit Metric : <tt class="ph tt">(counter_name).(rollup_metric)</tt></li>
<li class="li"> Sub-Metric : <tt class="ph tt">(counter_name).(rollup_metric).(submetric)</tt></li>
</ul>
where
<ul class="ul">
<li class="li"> unit: A logical or physical unit of the GPU
</li>
<li class="li"> subunit: The subunit within the unit where the counter was measured. Sometimes this is a pipeline mode instead.
</li>
<li class="li"> pipestage: The pipeline stage within the subunit where the counter was measured.
</li>
<li class="li"> quantity: What is being measured. Generally matches the <dfn class="term">dimensional units</dfn>.
</li>
<li class="li"> qualifiers: Any additional predicates or filters applied to the counter. Often, an unqualified counter can be broken down
into several qualified sub-components.
</li>
<li class="li"> interface: Of the form <tt class="ph tt">sender2receiver</tt>, where <tt class="ph tt">sender</tt> is the source-unit and <tt class="ph tt">receiver</tt> is the destination-unit.
</li>
<li class="li"> rollup_metric: One of sum, avg, min, max.
</li>
<li class="li"> submetric: refer to section <a class="xref" href="index.html#metrics-structure__metrics-entities" shape="rect">Metrics Entities</a></li>
</ul>
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.
</div>
</div>
<div class="section" id="metrics-structure__metrics-cycles"><a name="metrics-structure__metrics-cycles" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Cycle Metrics</h3>
<div class="p">
Counters using the term <tt class="ph tt">cycles</tt> in the name report the number of cycles in the unit's clock domain. Unit-level cycle metrics include:
<ul class="ul">
<li class="li"><tt class="ph tt">unit__cycles_elapsed</tt> : The number of cycles within a range. The cycles' DimUnits are specific to the unit's clock domain.
</li>
<li class="li"><tt class="ph tt">unit__cycles_active</tt> : The number of cycles where the unit was processing data.
</li>
<li class="li"><tt class="ph tt">unit__cycles_stalled</tt> : The number of cycles where the unit was unable to process new data because its output interface was blocked.
</li>
<li class="li"><tt class="ph tt">unit__cycles_idle</tt> : The number of cycles where the unit was idle.
</li>
</ul>
Interface-level cycle counters are often (not always) available in the following variations:
<ul class="ul">
<li class="li"><tt class="ph tt">unit__(interface)_active</tt> : Cycles where data was transferred from source-unit to destination-unit.
</li>
<li class="li"><tt class="ph tt">unit__(interface)_stalled</tt> : Cycles where the source-unit had data, but the destination-unit was unable to accept data.
</li>
</ul>
</div>
</div>
</div>
</div>
<div class="topic concept nested1" id="metrics-decoder"><a name="metrics-decoder" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#metrics-decoder" name="metrics-decoder" shape="rect">3.3. Metrics Decoder</a></h3>
<div class="body conbody">
<p class="p">
The following explains terms found in <span class="keyword">NVIDIA Nsight Compute</span> metric names, as introduced in <a class="xref" href="index.html#metrics-structure" shape="rect">Metrics Structure</a>.
</p>
<div class="section" id="metrics-decoder__metrics-hw-units"><a name="metrics-decoder__metrics-hw-units" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Units</h3>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">dram</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Device (main) memory, where the GPUs global and local memory resides.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">fbpa</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">fe</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The Frontend unit is responsible for the overall flow of workloads sent by the driver.
FE also facilitates a number of synchronization operations.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">gpc</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The General Processing Cluster contains SM, Texture and L1 in the form of TPC(s).
It is replicated several times across a chip.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">gpu</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The entire Graphics Processing Unit.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">gr</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Graphics Engine is responsible for all 2D and 3D graphics, compute work, and synchronous graphics copying work.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">idc</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The InDexed Constant Cache is a subunit of the SM responsible for caching constants that are indexed with a register.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">l1tex</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">lts</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sm</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sys</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Logical grouping of several units.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">tpc</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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).
</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-decoder__metrics-sub-units"><a name="metrics-decoder__metrics-sub-units" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Subunits</h3>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">aperture_device</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Memory interface to local device memory (dram)</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">aperture_peer</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Memory interface to remote device memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">aperture_sysmem</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Memory interface to system memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">global</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">lg</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Local/Global memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">local</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">lsu</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Load/Store unit</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">lsuin</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Load/Store input</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">mio</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Memory input/output</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">mioc</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Memory input/output control</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">shared</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">surface</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Surface memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">texin</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">TEXIN</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">texture</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Texture memory</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">xbar</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The Crossbar (XBAR) is responsible for carrying packets from a given source unit to a specific destination unit.</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-decoder__metrics-hw-pipelines"><a name="metrics-decoder__metrics-hw-pipelines" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Pipelines</h3>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">adu</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">alu</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">cbu</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Convergence Barrier Unit.
The CBU is responsible for warp-level convergence, barrier, and branch instructions.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">fma</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">fmaheavy</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Fused Multiply Add/Accumulate Heavy. FMAHeavy performs FP32 arithmetic (FADD, FMUL, FMAD), FP16 arithmetic (HADD2, HMUL2,
HFMA2), and integer dot products.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">fmalite</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Fused Multiply Add/Accumulate Lite. FMALite performs FP32 arithmetic (FADD, FMUL, FMA) and FP16 arithmetic (HADD2, HMUL2,
HFMA2).
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">fp16</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">fp64</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Double-precision floating-point. The implementation of FP64 varies greatly per chip.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">lsu</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">tex</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">tma</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Tensor Memory Access Unit.
Provides efficient data transfer mechanisms between global and shared memories with the ability to understand and traverse
multidimensional data layouts.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">uniform</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Uniform Data Path.
This scalar unit executes instructions where all threads use the same input and generate the same output.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">xu</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-decoder__metrics-quantities"><a name="metrics-decoder__metrics-quantities" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Quantities</h3>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">instruction</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">An assembly (SASS) instruction.
Each executed instruction may generate zero or more requests.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">request</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sector</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">tag</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">wavefront</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-decoder__metrics-details"><a name="metrics-decoder__metrics-details" shape="rect">
<!-- --></a><p class="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 <dfn class="term">request</dfn> 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 <dfn class="term">sectors</dfn> within these cache lines.
The L1TEX unit has internally multiple processing stages operating in a pipeline.
</p>
<p class="p">
A <dfn class="term">wavefront</dfn> 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 <dfn class="term">requests:sectors is 1:N, wavefronts:sectors 1:N, and requests:wavefronts is 1:N</dfn>.
</p>
<p class="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 <dfn class="term">access pattern</dfn> 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>
</div>
</div>
</div>
<div class="topic concept nested1" id="range-and-precision"><a name="range-and-precision" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#range-and-precision" name="range-and-precision" shape="rect">3.4. Range and Precision</a></h3>
<div class="body conbody">
<div class="section" id="range-and-precision__overview"><a name="range-and-precision__overview" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Overview</h3>
<p class="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>
</div>
<div class="section" id="range-and-precision__async-activity"><a name="range-and-precision__async-activity" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Asynchronous GPU activity</h3>
<p class="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>
</div>
<div class="section" id="range-and-precision__multi-pass-error"><a name="range-and-precision__multi-pass-error" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Multi-pass data collection</h3>
<p class="p">
Out-of-range metrics often occur when the profiler <a class="xref" href="index.html#kernel-replay" shape="rect">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 class="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>
</div>
<div class="section" id="range-and-precision__tool-issue"><a name="range-and-precision__tool-issue" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Tool issue</h3>
<p class="p">
If you still observe metric issues after following the guidelines above, please <a class="xref" href="https://forums.developer.nvidia.com/c/development-tools/nsight-compute" target="_blank" shape="rect">reach out to us</a> and describe your issue.
</p>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="metrics-reference"><a name="metrics-reference" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#metrics-reference" name="metrics-reference" shape="rect">4. Metrics Reference</a></h2>
<div class="body conbody">
<div class="section" id="metrics-reference__metrics-ref-overview"><a name="metrics-reference__metrics-ref-overview" shape="rect">
<!-- --></a><h2 class="title sectiontitle">Overview</h2>
<p class="p">
Most metrics in <span class="keyword">NVIDIA Nsight Compute</span> can be queried using the <span class="keyword">ncu</span> command
line interface's <a class="xref" href="../NsightComputeCli/index.html#command-line-options-profile" shape="rect">--query-metrics</a> option.
</p>
<p class="p">
The following metrics can be collected explicitly, but are not listed by <samp class="ph codeph">--query-metrics</samp>, and do not follow the
naming scheme explained in <a class="xref" href="index.html#metrics-structure" shape="rect">Metrics Structure</a>. They should be used as-is instead.
</p>
<p class="p"><samp class="ph codeph">launch__*</samp> 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="xref" href="../NsightCompute/index.html#occupancy-calculator" shape="rect">CUDA Occupancy Calculator</a>.
</p>
</div>
<div class="section" id="metrics-reference__metrics-ref-launch"><a name="metrics-reference__metrics-ref-launch" shape="rect">
<!-- --></a><h2 class="title sectiontitle">Launch Metrics</h2>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__block_dim_x</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of threads for the kernel launch in X dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__block_dim_y</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of threads for the kernel launch in Y dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__block_dim_z</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of threads for the kernel launch in Z dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__blocks_size</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Total number of threads per block for the kernel launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__cluster_dim_x</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of clusters for the kernel launch in X dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__cluster_dim_y</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of clusters for the kernel launch in Y dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__cluster_dim_z</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of clusters for the kernel launch in Z dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__cluster_max_active</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__cluster_max_potential_size</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Largest valid cluster size for the kernel function and launch configuration.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__cluster_scheduling_policy</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Cluster scheduling policy.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__context_id</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">CUDA context id for the kernel launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__device_id</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">CUDA device id for the kernel launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__func_cache_config</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__function_pcs</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Kernel function entry PCs.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__grid_dim_x</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of blocks for the kernel launch in X dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__grid_dim_y</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of blocks for the kernel launch in Y dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__grid_dim_z</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of blocks for the kernel launch in Z dimension.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__grid_size</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Total number of blocks for the kernel launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__occupancy_cluster_gpu_pct</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Overall GPU occupancy due to clusters.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__occupancy_cluster_pct</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The ratio of active blocks to the max possible active blocks due to clusters.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__occupancy_per_cluster_size</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of active clusters for given cluster size.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__registers_per_thread</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of registers allocated per thread.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__shared_mem_config_size</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__shared_mem_per_block_driver</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Shared memory size per block, allocated for the CUDA driver.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__shared_mem_per_block_dynamic</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Dynamic shared memory size per block, allocated for the kernel.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__shared_mem_per_block_static</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Static shared memory size per block, allocated for the kernel.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__stream_id</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">CUDA stream id for the kernel launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__thread_count</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Total number of threads across all blocks for the kernel launch.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">launch__waves_per_multiprocessor</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-reference__metrics-ref-nvlink-topology"><a name="metrics-reference__metrics-ref-nvlink-topology" shape="rect">
<!-- --></a><h2 class="title sectiontitle">NVLink Topology Metrics</h2>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__bandwidth</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Link bandwidth in bytes/s.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__count_logical</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Total number of logical NVLinks.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__count_physical</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Total number of physical links.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__destination_ports</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Destination port numbers (as strings).</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__dev0Id</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">ID of the first connected device.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__dev0type</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Type of the first connected device.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__dev1Id</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">ID of the second connected device.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__dev1type</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Type of the second connected device.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__dev_display_name_all</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Device display name.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__enabled_mask</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">NVLink enablement mask, per device.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__is_direct_link</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Indicates, per NVLink, if the link is direct.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__is_nvswitch_connected</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Indicates if NVSwitch is connected.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__max_count</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Maximum number of NVLinks.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__peer_access</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Indicates if peer access is supported.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__peer_atomic</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Indicates if peer atomics are supported.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__source_ports</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Source port numbers (as strings).</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__system_access</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Indicates if system access is supported.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">nvlink__system_atomic</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Indicates if system atomics are supported.</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-reference__metrics-ref-device"><a name="metrics-reference__metrics-ref-device" shape="rect">
<!-- --></a><h2 class="title sectiontitle">Device Attributes</h2>
<p class="p"><samp class="ph codeph">device__attribute_*</samp> metrics represent
<a class="xref" href="https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html#group__CUDART__TYPES_1g49e2f8c2c0bd6fe264f2fc970912e5cd" target="_blank" shape="rect">CUDA device attributes</a>.
Collecting them does not require an addition kernel replay pass, as their value is available from the CUDA driver for each
CUDA device.
</p>
</div>
<div class="section" id="metrics-reference__metrics-ref-source"><a name="metrics-reference__metrics-ref-source" shape="rect">
<!-- --></a><h2 class="title sectiontitle">Source Metrics</h2>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">branch_inst_executed</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of unique branch targets assigned to the instruction, including both divergent and uniform branches.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">derived__avg_thread_executed</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Average number of thread-level executed instructions per warp (regardless of their predicate). Computed as: thread_inst_executed
/ inst_executed
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">derived__avg_thread_executed_true</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Average number of predicated-on thread-level executed instructions per warp. Computed as: thread_inst_executed_true / inst_executed</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">derived__memory_l1_wavefronts_shared_excessive</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Excessive number of wavefronts in L1 from shared memory instructions, because not all not predicated-off threads performed
the operation.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">derived__memory_l2_theoretical_sectors_global_excessive</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Excessive theoretical number of sectors requested in L2 from global memory instructions, because not all not predicated-off
threads performed the operation.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">inst_executed</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_access_size_type</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The size of the memory access, in bits.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_access_type</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The type of memory access (e.g. load or store).</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_l1_tag_requests_global</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of L1 tag requests generated by global memory instructions.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_l1_wavefronts_shared</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of wavefronts in L1 from shared memory instructions.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_l1_wavefronts_shared_ideal</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Ideal number of wavefronts in L1 from shared memory instructions, assuming each not predicated-off thread performed the operation.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_l2_theoretical_sectors_global</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Theoretical number of sectors requested in L2 from global memory instructions.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_l2_theoretical_sectors_global_ideal</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Ideal number of sectors requested in L2 from global memory instructions, assuming each not predicated-off thread performed
the operation.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_l2_theoretical_sectors_local</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Theoretical number of sectors requested in L2 from local memory instructions.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">memory_type</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">The accessed address space (global/local/shared).</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__branch_targets_threads_divergent</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of divergent branch targets, including fallthrough. Incremented only when there are two or more active threads with
divergent targets.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__branch_targets_threads_uniform</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of uniform branch execution, including fallthrough, where all active threads selected the same branch target.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">thread_inst_executed</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of thread-level executed instructions, regardless of predicate presence or evaluation.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">thread_inst_executed_true</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of thread-level executed instructions, where the instruction predicate evaluated to true, or no predicate was given.</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-reference__metrics-ref-memdesc"><a name="metrics-reference__metrics-ref-memdesc" shape="rect">
<!-- --></a><h2 class="title sectiontitle">L2 Cache Eviction Metrics</h2>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__sass_inst_executed_memdesc_explicit_evict_type</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">L2 cache eviction policy types.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_first</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions with L2 cache eviction hit property 'first'.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_last</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions with L2 cache eviction hit property 'last'.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_normal</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions with L2 cache eviction hit property 'normal'.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__sass_inst_executed_memdesc_explicit_hitprop_evict_normal_demote</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions with L2 cache eviction hit property 'normal demote'.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__sass_inst_executed_memdesc_explicit_missprop_evict_first</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions with L2 cache eviction miss property 'first'.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">smsp__sass_inst_executed_memdesc_explicit_missprop_evict_normal</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions with L2 cache eviction miss property 'normal'.</td>
</tr>
</tbody>
</table>
</div>
</div>
<div class="section" id="metrics-reference__metrics-ref-sass-per-opcode"><a name="metrics-reference__metrics-ref-sass-per-opcode" shape="rect">
<!-- --></a><h2 class="title sectiontitle">Instructions Per Opcode Metrics</h2>
<div class="tablenoborder">
<table cellpadding="4" cellspacing="0" summary="" class="table" frame="border" border="1" rules="all">
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sass__inst_executed_per_opcode</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions, instanced by basic SASS opcode.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sass__inst_executed_per_opcode_with_modifier_all</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions, instanced by all SASS opcode modifiers.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sass__inst_executed_per_opcode_with_modifier_selective</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of warp-level executed instructions, instanced by selective SASS opcode modifiers.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sass__thread_inst_executed_true_per_opcode</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of thread-level executed instructions, instanced by basic SASS opcode.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sass__thread_inst_executed_true_per_opcode_with_modifier_all</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of thread-level executed instructions, instanced by all SASS opcode modifiers.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" rowspan="1" colspan="1"><tt class="ph tt">sass__thread_inst_executed_true_per_opcode_with_modifier_selective</tt></td>
<td class="entry" valign="top" rowspan="1" colspan="1">Number of thread-level executed instructions, instanced by selective SASS opcode modifiers.</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="sampling"><a name="sampling" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#sampling" name="sampling" shape="rect">5. Sampling</a></h2>
<div class="body conbody">
<p class="p"><span class="keyword">NVIDIA Nsight Compute</span> supports periodic sampling of the warp program counter and warp scheduler state on desktop devices of compute capability
6.1 and above.
</p>
<p class="p">
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>
</div>
<div class="topic concept nested1" id="statistical-sampler"><a name="statistical-sampler" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#statistical-sampler" name="statistical-sampler" shape="rect">5.1. Warp Scheduler States</a></h3>
<div class="body conbody">
<div class="p">
<div class="tablenoborder"><a name="statistical-sampler__warp-scheduler-states" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="statistical-sampler__warp-scheduler-states" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 2. Warp Scheduler States</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="20%" id="d54e3775" rowspan="1" colspan="1">State</th>
<th class="entry" valign="top" width="20%" id="d54e3778" rowspan="1" colspan="1">Hardware Support</th>
<th class="entry" valign="top" width="60%" id="d54e3781" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Allocation</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2-6.1</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">Warp was stalled waiting for a branch to resolve, waiting for all memory operations to retire, or waiting to be allocated
to the micro-scheduler.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Barrier</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
Also, try to identify which barrier instruction causes the most stalls, and optimize the code executed before that synchronization
point first.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Branch Resolving</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">7.0+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
Warp was stalled waiting for a branch target to be computed, and the warp program counter to be updated.
Excessively jumping (branching) can lead to more warps stalled for this reason.
See also the related <dfn class="term">No Instructions</dfn> state.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Dispatch</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Drain</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
Warp was stalled after EXIT waiting for all memory instructions 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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">IMC Miss</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">LG Throttle</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">7.0+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
If applicable, consider combining multiple lower-width memory operations into fewer wider memory operations
and try interleaving memory operations and math instructions.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Long Scoreboard</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
Warp was stalled waiting for a scoreboard dependency on a L1TEX (local, global, surface, tex) operation.
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,
or by changing the cache configuration, and consider moving frequently used data to shared memory.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Math Pipe Throttle</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Membar</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">MIO Throttle</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Misc</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">Warp was stalled for a miscellaneous hardware reason.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">No Instructions</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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 <dfn class="term">Branch Resolving</dfn> state.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Not Selected</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Selected</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">Warp was selected by the micro scheduler and issued an instruction.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Short Scoreboard</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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).
Verify if there are shared memory operations and reduce bank conflicts, if applicable.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Sleeping</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">7.0+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Tex Throttle</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">
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.
If applicable, consider combining multiple lower-width memory operations into fewer wider memory operations
and try interleaving memory operations and math instructions.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e3775" rowspan="1" colspan="1">Wait</td>
<td class="entry" valign="top" width="20%" headers="d54e3778" rowspan="1" colspan="1">5.2+</td>
<td class="entry" valign="top" width="60%" headers="d54e3781" rowspan="1" colspan="1">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.
If possible, try to further increase the number of active warps to hide the corresponding instruction latencies.
</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="reproducibility"><a name="reproducibility" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#reproducibility" name="reproducibility" shape="rect">6. Reproducibility</a></h2>
<div class="body conbody">
<p class="p">
In order to provide actionable and deterministic results across application runs,
<span class="keyword">NVIDIA Nsight Compute</span> applies various methods to adjust how metrics are collected.
This includes <a class="xref" href="index.html#serialization" shape="rect">serializing</a> kernel launches,
<a class="xref" href="index.html#cache-control" shape="rect">purging GPU caches</a> before each kernel replay
or <a class="xref" href="index.html#clock-control" shape="rect">adjusting GPU clocks</a>.
</p>
</div>
<div class="topic concept nested1" id="serialization"><a name="serialization" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#serialization" name="serialization" shape="rect">6.1. Serialization</a></h3>
<div class="body conbody">
<p class="p"><span class="keyword">NVIDIA Nsight Compute</span> 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 class="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 <samp class="ph codeph">TMPDIR/nsight-compute-lock</samp> is used.
On Windows, <dfn class="term">TMPDIR</dfn> is the path returned by the Windows <samp class="ph codeph">GetTempPath</samp> API function.
On other platforms, it is the path supplied by the first environment variable in the list
<samp class="ph codeph">TMPDIR, TMP, TEMP, TEMPDIR</samp>. If none of these is found, it's <samp class="ph codeph">/var/nvidia</samp> on QNX and <samp class="ph codeph">/tmp</samp> otherwise.
</p>
<p class="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 class="p">
It is currently not possible to disable this tool behavior.
Refer to the <a class="xref" href="index.html#faq" shape="rect">FAQ</a> entry on possible workarounds.
</p>
</div>
</div>
<div class="topic concept nested1" id="clock-control"><a name="clock-control" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#clock-control" name="clock-control" shape="rect">6.2. Clock Control</a></h3>
<div class="body conbody">
<p class="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 class="p">
To mitigate this non-determinism, <span class="keyword">NVIDIA Nsight Compute</span> attempts to limit GPU clock frequencies to their <dfn class="term">base</dfn> 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 class="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 <samp class="ph codeph">--clock-control</samp> option to specify if any clock frequencies should be fixed by the tool.
</p>
<div class="p">
Factors affecting Clock Control:
<ul class="ul">
<li class="li">Note that thermal throttling directed by the driver cannot be controlled by the tool and always overrides any selected options.</li>
<li class="li">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 <samp class="ph codeph">--clock-control</samp> to set the GPU clocks will fail or will be silently ignored when profiling on a GPU partition.
<ul class="ul">
<li class="li">On L4T, you can use the jetson_clocks script to lock the clocks at their maximums during profiling. </li>
</ul>
</li>
<li class="li">See the <a class="xref" href="index.html#special-configurations" shape="rect">Special Configurations</a> section for MIG and vGPU clock control.
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested1" id="cache-control"><a name="cache-control" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#cache-control" name="cache-control" shape="rect">6.3. Cache Control</a></h3>
<div class="body conbody">
<p class="p">
As explained in <a class="xref" href="index.html#kernel-replay" shape="rect">Kernel Replay</a>, the kernel might need to be replayed multiple times to collect all requested metrics.
While <span class="keyword">NVIDIA Nsight Compute</span> 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 class="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 class="p">
In order to make HW performance counter value more deterministic, <span class="keyword">NVIDIA Nsight Compute</span> 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 class="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 <samp class="ph codeph">--cache-control none</samp> to disable flushing of any HW cache by the tool.
</p>
</div>
</div>
<div class="topic concept nested1" id="persistence-mode"><a name="persistence-mode" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#persistence-mode" name="persistence-mode" shape="rect">6.4. Persistence Mode</a></h3>
<div class="body conbody">
<p class="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 class="p">
If <a class="xref" href="https://docs.nvidia.com/deploy/driver-persistence/index.html" target="_blank" shape="rect">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 class="p">
It is recommended to enable persistence mode on applicable operating systems before profiling with <span class="keyword">NVIDIA Nsight Compute</span> for more consistent application behavior.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="special-configurations"><a name="special-configurations" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#special-configurations" name="special-configurations" shape="rect">7. Special Configurations</a></h2>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested1" id="special-configurations-mig"><a name="special-configurations-mig" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#special-configurations-mig" name="special-configurations-mig" shape="rect">7.1. Multi Instance GPU</a></h3>
<div class="body conbody">
<p class="p"><span class="keyword">Multi-Instance GPU (MIG)</span> 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 <span class="keyword">GPU Instance</span>s. Each <span class="keyword">GPU Instance</span> 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 <span class="keyword">GPU Instance</span> can be further partitioned into one or more <span class="keyword">Compute Instance</span>s. Each <span class="keyword">Compute Instance</span> has exclusive ownership of its assigned SMs of the <span class="keyword">GPU Instance</span>. However, all <span class="keyword">Compute Instance</span>s within a <span class="keyword">GPU Instance</span> share the <span class="keyword">GPU Instance</span>'s memory and memory bandwidth. Every <span class="keyword">Compute Instance</span> acts and operates as a CUDA device with a unique device ID.
See the driver release notes as well as the documentation for the <samp class="ph codeph">nvidia-smi</samp> CLI tool for more information on how to configure MIG instances.
</p>
<p class="p">
For profiling, a <span class="keyword">Compute Instance</span> can be of one of two types: <dfn class="term">isolated</dfn> or <dfn class="term">shared</dfn>.
</p>
<p class="p">
An <dfn class="term">isolated</dfn> <span class="keyword">Compute Instance</span> owns all of its assigned resources and does not share any GPU unit with another <span class="keyword">Compute Instance</span>. In other words, the <span class="keyword">Compute Instance</span> is the same size as its parent <span class="keyword">GPU Instance</span> and consequently does not have any other sibling <span class="keyword">Compute Instance</span>s. Profiling works as usual for isolated <span class="keyword">Compute Instance</span>s.
</p>
<p class="p">
A <dfn class="term">shared</dfn> <span class="keyword">Compute Instance</span> uses GPU resources that can potentially also be accessed by other <span class="keyword">Compute Instance</span>s in the same <span class="keyword">GPU Instance</span>. 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 <samp class="ph codeph">==ERROR== Failed to access the following metrics. When profiling on a MIG instance, it is not possible to collect metrics
from GPU units that are shared with other MIG instances</samp> followed by the list of failing metrics. Collecting only metrics from GPU units that are exclusively owned by a shared <span class="keyword">Compute Instance</span> is still possible.
</p>
<div class="section" id="special-configurations-mig__clocks"><a name="special-configurations-mig__clocks" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Locking Clocks</h3>
<p class="p"><span class="keyword">NVIDIA Nsight Compute</span> is not able to set the clock frequency on any <span class="keyword">Compute Instance</span> for profiling.
You can continue analyzing kernels without fixed clock frequencies (using <samp class="ph codeph">--clock-control none</samp>; see <a class="xref" href="index.html#clock-control" shape="rect">here</a> for more details).
If you have sufficient permissions, <samp class="ph codeph">nvidia-smi</samp> can be used to configure a fixed frequency for the whole GPU by calling <samp class="ph codeph">nvidia-smi --lock-gpu-clocks=tdp,tdp</samp>.
This sets the GPU clocks to the base TDP frequency until you reset the clocks by calling <samp class="ph codeph">nvidia-smi --reset-gpu-clocks</samp>.
</p>
</div>
<div class="section" id="special-configurations-mig__baremetal"><a name="special-configurations-mig__baremetal" shape="rect">
<!-- --></a><h3 class="title sectiontitle">MIG on Baremetal (non-vGPU)</h3>
<p class="p">
All <span class="keyword">Compute Instance</span>s on a GPU share the same clock frequencies.
</p>
</div>
<div class="section" id="special-configurations-mig__vgpu"><a name="special-configurations-mig__vgpu" shape="rect">
<!-- --></a><h3 class="title sectiontitle">MIG on NVIDIA vGPU</h3>
<p class="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 <span class="keyword">Compute Instance</span>s.
</p>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="roofline"><a name="roofline" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#roofline" name="roofline" shape="rect">8. Roofline Charts</a></h2>
<div class="body conbody">
<p class="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>
</div>
<div class="topic concept nested1" id="roofline-overview"><a name="roofline-overview" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#roofline-overview" name="roofline-overview" shape="rect">8.1. Overview</a></h3>
<div class="body conbody">
<p class="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
<dfn class="term">Arithmetic Intensity</dfn> (a ratio between <dfn class="term">Work</dfn> and <dfn class="term">Memory Traffic</dfn>), 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>
<div class="fig fignone"><span class="desc figdesc">Roofline overview.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/roofline-overview.png"></img></div><br clear="none"></br></div>
<div class="p">
This chart actually shows two different rooflines. However, the following components can
be identified for each:
<ul class="ul">
<li class="li"><strong class="ph b"><dfn class="term">Vertical Axis</dfn></strong> - The vertical axis represents <dfn class="term">Floating Point Operations per Second</dfn> (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.
</li>
<li class="li"><strong class="ph b"><dfn class="term">Horizontal Axis</dfn></strong> - The horizontal axis represents <dfn class="term">Arithmetic Intensity</dfn>, which is the ratio
between <dfn class="term">Work</dfn> (expressed in floating point operations per second), and <dfn class="term">Memory Traffic</dfn> (expressed
in bytes per second). The resulting unit is in floating point operations per byte. This axis is also shown using a logarithmic
scale.
</li>
<li class="li"><strong class="ph b"><dfn class="term">Memory Bandwidth Boundary</dfn></strong> - The memory bandwidth boundary is the <em class="ph i">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 <dfn class="term">SpeedOfLight_RooflineChart.section</dfn> file if desired.
</li>
<li class="li"><strong class="ph b"><dfn class="term">Peak Performance Boundary</dfn></strong> - The peak performance boundary is the <em class="ph i">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 <dfn class="term">SpeedOfLight_RooflineChart.section</dfn> file if desired.
</li>
<li class="li"><strong class="ph b"><dfn class="term">Ridge Point</dfn></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.
</li>
<li class="li"><strong class="ph b"><dfn class="term">Achieved Value</dfn></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.
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested1" id="roofline-analysis"><a name="roofline-analysis" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#roofline-analysis" name="roofline-analysis" shape="rect">8.2. Analysis</a></h3>
<div class="body conbody">
<p class="p">
The roofline chart can be very helpful in guiding performance optimization efforts for a particular kernel.
</p>
<div class="fig fignone"><span class="desc figdesc">Roofline anaysis.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/roofline-analysis.png"></img></div><br clear="none"></br></div>
<p class="p">
As shown here, the <dfn class="term">ridge point</dfn> partitions the roofline chart into two regions. The area shaded
in blue under the sloped <dfn class="term">Memory Bandwidth Boundary</dfn> is the <dfn class="term">Memory Bound</dfn> region, while the
area shaded in green under the <dfn class="term">Peak Performance Boundary</dfn> is the <dfn class="term">Compute Bound</dfn> region.
The region in which the <dfn class="term">achieved value</dfn> falls, determines the current limiting factor of kernel performance.
</p>
<p class="p">
The distance from the <dfn class="term">achieved value</dfn> to the respective roofline boundary (shown in this figure as a dotted
white line), represents the opportunity for performance improvement. The closer the <dfn class="term">achieved value</dfn> is to
the roofline boundary, the more optimal is its performance. An <dfn class="term">achieved value</dfn> that lies on the
<dfn class="term">Memory Bandwidth Boundary</dfn> but is not yet at the height of the <dfn class="term">ridge point</dfn> would indicate that
any further improvements in overall FLOP/s are only possible if the <dfn class="term">Arithmetic Intensity</dfn> is increased at
the same time.
</p>
<p class="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>
</div>
</div>
</div>
<div class="topic concept nested0" id="memory-chart"><a name="memory-chart" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#memory-chart" name="memory-chart" shape="rect">9. Memory Chart</a></h2>
<div class="body conbody">
<p class="p">
The <dfn class="term">Memory Chart</dfn> 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>
</div>
<div class="topic concept nested1" id="memory-chart-overview"><a name="memory-chart-overview" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-chart-overview" name="memory-chart-overview" shape="rect">9.1. Overview</a></h3>
<div class="body conbody">
<div class="fig fignone"><span class="desc figdesc">Memory chart for an NVIDIA A100 GPU</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/memory-chart-a100.png"></img></div><br clear="none"></br></div>
<div class="section" id="memory-chart-overview__memory-chart-logical-units"><a name="memory-chart-overview__memory-chart-logical-units" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Logical Units (green)</h3>
Logical units are shown in green color.
<ul class="ul">
<li class="li">Kernel: The CUDA kernel executing on the GPU's Streaming Multiprocessors</li>
<li class="li">Global: CUDA <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">global memory</a></li>
<li class="li">Local: CUDA <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">local memory</a></li>
<li class="li">Texture: CUDA <a class="xref" href="index.html#metrics-hw-model__metrics-hw-tex-surf" shape="rect">texture memory</a></li>
<li class="li">Surface: CUDA <a class="xref" href="index.html#metrics-hw-model__metrics-hw-tex-surf" shape="rect">surface memory</a></li>
<li class="li">Shared: CUDA <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">shared memory</a></li>
<li class="li">Load Global Store Shared: Instructions loading directly from global into shared memory without intermediate register file
access
</li>
</ul>
</div>
<div class="section" id="memory-chart-overview__memory-chart-physical-units"><a name="memory-chart-overview__memory-chart-physical-units" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Physical Units (blue)</h3>
Physical units are shown in blue color.
<ul class="ul">
<li class="li">
L1/TEX Cache: The <a class="xref" href="index.html#metrics-hw-model__metrics-hw-caches" shape="rect">L1/Texture cache</a>.
The underlying physical memory is split between this cache and the user-managed <dfn class="term">Shared Memory</dfn>.
</li>
<li class="li">
Shared Memory: CUDA's user-managed <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">shared memory</a>.
The underlying physical memory is split between this and the <dfn class="term">L1/TEX Cache</dfn>.
</li>
<li class="li">L2 Cache: The <a class="xref" href="index.html#metrics-hw-model__metrics-hw-caches" shape="rect">L2 cache</a></li>
<li class="li">L2 Compression: The memory compression unit of the <dfn class="term">L2 Cache</dfn></li>
<li class="li">System Memory: Off-chip <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">system (CPU) memory</a></li>
<li class="li">Device Memory: On-chip <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">device (GPU) memory</a> of the CUDA device that executes the kernel
</li>
<li class="li">Peer Memory: On-chip <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">device (GPU) memory</a> of other CUDA devices
</li>
</ul>
<p class="p">
Depending on the exact GPU architecture, the exact set of shown units can vary, as not all GPUs have all units.
</p>
</div>
<div class="section" id="memory-chart-overview__memory-chart-links"><a name="memory-chart-overview__memory-chart-links" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Links</h3>
<p class="p">
Links between <dfn class="term">Kernel</dfn> and other logical units represent the number of executed instructions (<dfn class="term">Inst</dfn>) targeting the respective unit.
For example, the link between <dfn class="term">Kernel</dfn> and <dfn class="term">Global</dfn> represents the instructions loading from or storing to the global memory space.
Instructions using the NVIDIA A100's <dfn class="term">Load Global Store Shared</dfn> paradigm are shown separately, as their register or cache access behavior
can be different from regular global loads or shared stores.
</p>
<p class="p">
Links between logical units and blue, physical units represent the number of requests (<dfn class="term">Req</dfn>) issued as a result of their respective instructions.
For example, the link going from <dfn class="term">L1/TEX Cache</dfn> to <dfn class="term">Global</dfn> shows the number of requests generated due to global load instructions.
</p>
<p class="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 class="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 class="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>
<div class="fig fignone"><span class="desc figdesc">Mapping of peak values between memory tables and memory chart</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/memory-peak-mapping.png"></img></div><br clear="none"></br></div>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="memory-tables"><a name="memory-tables" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#memory-tables" name="memory-tables" shape="rect">10. Memory Tables</a></h2>
<div class="body conbody">
<p class="p">
The <dfn class="term">Memory Tables</dfn> 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 <dfn class="term">UNUSED</dfn> in the tooltip.
You can hover over row or column headers to see a description of this part of the table.
</p>
</div>
<div class="topic concept nested1" id="memory-tables-smem"><a name="memory-tables-smem" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-tables-smem" name="memory-tables-smem" shape="rect">10.1. Shared Memory</a></h3>
<div class="body conbody">
<div class="fig fignone"><span class="desc figdesc">Example Shared Memory table, collected on an RTX 2080 Ti</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/memory-tables-smem.png"></img></div><br clear="none"></br></div>
<div class="section" id="memory-tables-smem__memory-tables-smem-columns"><a name="memory-tables-smem__memory-tables-smem-columns" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Columns</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Instructions</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
<p class="p">
For each access type, the total number of all actually executed assembly (SASS) <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">instructions</a> per warp.
Predicated-off instructions are not included.
</p>
<p class="p">
E.g., the instruction <dfn class="term">STS</dfn> would be counted towards <dfn class="term">Shared Store</dfn>.
</p>
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Requests</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
<p class="p">
The total number of all <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">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="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Wavefronts</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Number of <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">wavefronts</a> required to service the requested shared memory data.
Wavefronts are serialized and processed on different cycles.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">% Peak</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Bank Conflicts</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
</table>
</div>
<div class="section" id="memory-tables-smem__memory-tables-smem-rows"><a name="memory-tables-smem__memory-tables-smem-rows" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Rows</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">(Access Types)</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Shared memory access operations.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The aggregate for all access types in the same column.
</td>
</tr>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="memory-tables-l1"><a name="memory-tables-l1" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-tables-l1" name="memory-tables-l1" shape="rect">10.2. L1/TEX Cache</a></h3>
<div class="body conbody">
<div class="fig fignone"><span class="desc figdesc">Example L1/TEX Cache memory table, collected on an RTX 2080 Ti</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/memory-tables-l1.png"></img></div><br clear="none"></br></div>
<div class="fig fignone"><span class="desc figdesc">Model of the Global Load Pipeline for the L1TEX cache on GA100, mapped to the memory table.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/hw-model-l1tex-ga100-global.png"></img></div><br clear="none"></br></div>
<div class="section" id="memory-tables-l1__memory-tables-l1-columns"><a name="memory-tables-l1__memory-tables-l1-columns" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Columns</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Instructions</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
<p class="p">
For each access type, the total number of all actually executed assembly (SASS) <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">instructions</a> per warp.
Predicated-off instructions are not included.
</p>
<p class="p">
E.g., the instruction <dfn class="term">LDG</dfn> would be counted towards <dfn class="term">Global Loads</dfn>.
</p>
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Requests</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
<p class="p">
The total number of all <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">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 class="p">
In the example, each of the 65536 global load instructions generates exactly one request.
</p>
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Wavefronts</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Number of <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">wavefronts</a> required to service the requested memory operation.
Wavefronts are serialized and processed on different cycles.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Wavefront % Peak</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Percentage of peak utilization for the units processing <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">wavefronts</a>.
High numbers can imply that the processing pipelines are saturated and can become a bottleneck.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sectors</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The total number of all L1 <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sectors/Req</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
<p class="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="xref" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory-accesses" target="_blank" shape="rect">uncoalesced memory accesses</a>
and will result in increased memory traffic.
</p>
<p class="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="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Hit Rate</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1"><a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">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 <dfn class="term">Sector Misses to L2</dfn>.
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 <dfn class="term">Tag Hit Rate</dfn> (not shown).
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Bytes</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sector Misses to L2</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
<p class="p">
Total number of sectors that miss in L1 and generate subsequent requests in the <a class="xref" href="index.html#memory-tables-l2" shape="rect">L2 Cache</a>.
</p>
<p class="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="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">% Peak to L2</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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="xref" href="index.html#statistical-sampler__warp-scheduler-states" shape="rect">warp stalls</a>.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Returns to SM</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Number of return packets sent from the L1 cache back to the SM.
Larger request access sizes result in higher number of returned packets.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">% Peak to SM</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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="xref" href="index.html#statistical-sampler__warp-scheduler-states" shape="rect">warp stalls</a>.
Improving read-coalescing or the <dfn class="term">L1 hit rate</dfn> could reduce this utilization.
</td>
</tr>
</table>
</div>
<div class="section" id="memory-tables-l1__memory-tables-l1-rows"><a name="memory-tables-l1__memory-tables-l1-rows" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Rows</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">(Access Types)</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The various access types, e.g. loads from global memory or reduction operations on surface memory.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Loads</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The aggregate of all load access types in the same column.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Stores</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The aggregate of all store access types in the same column.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The aggregate of all load and store access types in the same column.
</td>
</tr>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="memory-tables-l2"><a name="memory-tables-l2" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-tables-l2" name="memory-tables-l2" shape="rect">10.3. L2 Cache</a></h3>
<div class="body conbody">
<div class="fig fignone"><span class="desc figdesc">Example L2 Cache memory table, collected on an RTX 2080 Ti</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/memory-tables-l2.png"></img></div><br clear="none"></br></div>
<div class="fig fignone"><span class="desc figdesc">Model of the L2 cache on GA100, mapped to the memory table.</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/hw-model-lts-ga100.png"></img></div><br clear="none"></br></div>
<div class="section" id="memory-tables-l2__memory-tables-l2-columns"><a name="memory-tables-l2__memory-tables-l2-columns" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Columns</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Requests</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
For each access type, the total number of <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">requests</a> made to the L2 cache.
This correlates with the <a class="xref" href="index.html#memory-tables-l1__memory-tables-l1-columns" shape="rect">Sector Misses to L2</a> for the L1 cache.
Each request targets one 128 byte cache line.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sectors</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
For each access type, the total number of <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">sectors</a> requested from the L2 cache.
Each request accesses one or more sectors.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sectors/Req</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
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="xref" href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory-accesses" target="_blank" shape="rect">uncoalesced memory accesses</a>
and will result in increased memory traffic.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">% Peak</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Hit Rate</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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 <dfn class="term">Sector Misses to Device</dfn>, <dfn class="term">Sector Misses to System</dfn>, or <dfn class="term">Sector Misses to Peer</dfn>.
Higher hit rates imply better performance due to lower access latencies, as the request can be served by L2 instead of a later
stage.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Bytes</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Throughput</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Achieved L2 cache throughput in bytes per second.
High values indicate high utilization of the unit.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sector Misses to Device</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total number of sectors that miss in L2 and generate <a class="xref" href="index.html#memory-tables-dram" shape="rect">subsequent requests</a>
in <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">device memory</a>.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sector Misses to System</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total number of sectors that miss in L2 and generate subsequent requests
in <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">system memory</a>.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sector Misses to Peer</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total number of sectors that miss in L2 and generate subsequent requests
in <a class="xref" href="index.html#metrics-hw-model__metrics-hw-memory" shape="rect">peer memory</a>.
</td>
</tr>
</table>
</div>
<div class="section" id="memory-tables-l2__memory-tables-l2-rows"><a name="memory-tables-l2__memory-tables-l2-rows" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Rows</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">(Access Types)</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The various access types, e.g. loads or reductions originating from L1 cache.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">L1/TEX Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total for all operations originating from the L1 cache.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">ECC Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">L2 Fabric Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">GPU Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total for all operations across all clients of the L2 cache.
Independent of having them split out separately in this table.
</td>
</tr>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="memory-tables-l2-evict-policy"><a name="memory-tables-l2-evict-policy" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-tables-l2-evict-policy" name="memory-tables-l2-evict-policy" shape="rect">10.4. L2 Cache Eviction Policies</a></h3>
<div class="body conbody">
<div class="fig fignone"><span class="desc figdesc">Example L2 Cache Eviction Policies memory table, collected on an A100 GPU</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/memory-tables-l2-evict-policy.png"></img></div><br clear="none"></br></div>
<div class="section" id="memory-tables-l2-evict-policy__memory-tables-l2-evict-policy-columns"><a name="memory-tables-l2-evict-policy__memory-tables-l2-evict-policy-columns" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Columns</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">First</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Number of sectors accessed in the L2 cache using the <samp class="ph codeph">evict_first</samp> 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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Hit Rate</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Cache hit rate for sector accesses in the L2 cache using the <samp class="ph codeph">evict_first</samp> policy.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Last</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Number of sectors accessed in the L2 cache using the <samp class="ph codeph">evict_last</samp> policy.
Data cached with this policy will be last in the eviction priority order and will likely be evicted only after other data
with <samp class="ph codeph">evict_normal</samp> or <samp class="ph codeph">evict_first</samp> eviction policy is already evicted.
This policy is suitable for data that should remain persistent in cache.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Hit Rate</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Cache hit rate for sector accesses in the L2 cache using the <samp class="ph codeph">evict_last</samp> policy.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Normal</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Number of sectors accessed in the L2 cache using the <samp class="ph codeph">evict_normal</samp> policy.
This is the default policy.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Hit Rate</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Cache hit rate for sector accesses in the L2 cache using the <samp class="ph codeph">evict_normal</samp> policy.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Normal Demote</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Number of sectors accessed in the L2 cache using the <samp class="ph codeph">evict_normal_demote</samp> policy.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Hit Rate</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Cache hit rate for sector accesses in the L2 cache using the <samp class="ph codeph">evict_normal_demote</samp> policy.
</td>
</tr>
</table>
</div>
<div class="section" id="memory-tables-l2-evict-policy__memory-tables-l2-evict-policy-rows"><a name="memory-tables-l2-evict-policy__memory-tables-l2-evict-policy-rows" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Rows</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">(Access Types)</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The various access types, e.g. loads or reductions, originating from L1 cache.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">L1/TEX Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total for all operations originating from the L1 cache.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">L2 Fabric Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">GPU Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total for all operations across all clients of the L2 cache.
Independent of having them split out separately in this table.
</td>
</tr>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="memory-tables-dram"><a name="memory-tables-dram" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memory-tables-dram" name="memory-tables-dram" shape="rect">10.5. Device Memory</a></h3>
<div class="body conbody">
<div class="fig fignone"><span class="desc figdesc">Example Device Memory table, collected on an RTX 2080 Ti</span><br clear="none"></br><div class="imagecenter"><img class="image imagecenter" src="graphics/memory-tables-dram.png"></img></div><br clear="none"></br></div>
<div class="section" id="memory-tables-dram__memory-tables-dram-columns"><a name="memory-tables-dram__memory-tables-dram-columns" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Columns</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Sectors</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
For each access type, the total number of <a class="xref" href="index.html#metrics-decoder__metrics-quantities" shape="rect">sectors</a> requested from device memory.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">% Peak</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
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.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Bytes</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Total number of bytes transferred between <a class="xref" href="index.html#memory-tables-l2" shape="rect">L2 Cache</a> and device memory.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Throughput</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Achieved device memory throughput in bytes per second.
High values indicate high utilization of the unit.
</td>
</tr>
</table>
</div>
<div class="section" id="memory-tables-dram__memory-tables-dram-rows"><a name="memory-tables-dram__memory-tables-dram-rows" shape="rect">
<!-- --></a><h3 class="title sectiontitle">Rows</h3>
<table cellpadding="4" cellspacing="0" summary="" border="1" class="simpletable">
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">(Access Types)</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
Device memory loads and stores.
</td>
</tr>
<tr class="strow">
<td valign="top" class="stentry" rowspan="1" colspan="1"><tt class="ph tt">Total</tt></td>
<td valign="top" class="stentry" rowspan="1" colspan="1">
The aggregate for all access types in the same column.
</td>
</tr>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="faq"><a name="faq" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#faq" name="faq" shape="rect">11. FAQ</a></h2>
<div class="body conbody">
<ul class="ul">
<li class="li"><strong class="ph b">n/a metric values</strong><p class="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="xref" href="index.html#metrics-structure__metrics-examples" shape="rect">suffix</a>.
Verify the metric name against the output of of the <samp class="ph codeph">--query-metrics</samp><span class="keyword">NVIDIA Nsight Compute CLI</span> option.
</p>
<p class="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 class="p">
Finally, the metric might simply not exist for the targeted GPU architecture. For example, the IMMA pipeline metric
<samp class="ph codeph">sm__inst_executed_pipe_tensor_op_imma.avg.pct_of_peak_sustained_active</samp> is not available on GV100 chips.
</p>
</li>
<li class="li"><strong class="ph b">Metric values outside the expected logical range</strong><p class="p">
This includes e.g. percentages exceeding 100% or metrics reporting negative values.
For further details, see <a class="xref" href="index.html#range-and-precision" shape="rect">Range and Precision</a>.
</p>
</li>
<li class="li"><strong class="ph b">ERR_NVGPUCTRPERM - The user does not have permission to access NVIDIA GPU Performance Counters on the target device.</strong><p class="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="xref" href="https://developer.nvidia.com/ERR_NVGPUCTRPERM" target="_blank" shape="rect">https://developer.nvidia.com/ERR_NVGPUCTRPERM</a>.
</p>
</li>
<li class="li"><strong class="ph b">Unsupported GPU</strong><p class="p">
This indicates that the GPU, on which the current kernel is launched, is not supported.
See the <a class="xref" href="../ReleaseNotes/index.html#gpu-support" shape="rect">Release Notes</a>
for a list of devices supported by your version of <span class="keyword">NVIDIA Nsight Compute</span>.
It can also indicate that the current <dfn class="term">GPU configuration</dfn> is not supported.
For example, <span class="keyword">NVIDIA Nsight Compute</span> might not be able to profile GPUs in SLI configuration.
</p>
</li>
<li class="li"><strong class="ph b">Connection error detected communicating with target application.</strong><p class="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 class="li"><strong class="ph b">Failed to connect. The target process may have exited.</strong><div class="p">
This occurs if
<ul class="ul">
<li class="li">the application does not call any CUDA API calls before it exits.</li>
<li class="li">
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 <dfn class="term">Connection Dialog</dfn>.
</li>
<li class="li">the application crashes before calling any CUDA API calls.</li>
<li class="li">the application launches child processes which use the CUDA. In this case, launch with the <samp class="ph codeph">--target-processes all</samp> option.
</li>
</ul>
</div>
</li>
<li class="li"><strong class="ph b">The profiler returned an error code: (number)</strong><p class="p">
For the non-interactive <dfn class="term">Profile</dfn> activity, the <span class="keyword">NVIDIA Nsight Compute CLI</span> is started to generate the report.
If either the application exited with a non-zero return code, or the <span class="keyword">NVIDIA Nsight Compute CLI</span> encountered an error itself,
the resulting return code will be shown in this message.
</p>
<p class="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 class="p">
To debug this issue, it can help to run the data collection directly from the command line using <samp class="ph codeph"><span class="keyword">ncu</span></samp>
in order to observe the application's and the profiler's command line output, e.g. <samp class="ph codeph">==ERROR== The application returned an error code (11)</samp></p>
</li>
<li class="li"><strong class="ph b">Failed to open/create lock file (path). Please check that this process has write permissions on this file.</strong><p class="p"><span class="keyword">NVIDIA Nsight Compute</span> failed to create or open the file <samp class="ph codeph">(path)</samp> with write permissions.
This file is used for inter-process <a class="xref" href="index.html#serialization" shape="rect">serialization</a>.
<span class="keyword">NVIDIA Nsight Compute</span> 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 class="p">
The file is in the current temporary directory, i.e. <samp class="ph codeph">TMPDIR/nsight-compute-lock</samp>.
On Windows, <dfn class="term">TMPDIR</dfn> is the path returned by the Windows <samp class="ph codeph">GetTempPath</samp> API function.
On other platforms, it is the path supplied by the first environment variable in the list
<samp class="ph codeph">TMPDIR, TMP, TEMP, TEMPDIR</samp>. If none of these is found, it's <samp class="ph codeph">/var/nvidia</samp> on QNX and <samp class="ph codeph">/tmp</samp> otherwise.
</p>
<p class="p">
Older versions of <span class="keyword">NVIDIA Nsight Compute</span> 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>
<div class="p">
The following workarounds can be used to solve this problem:
<ul class="ul">
<li class="li">
If it is otherwise ensured that no concurrent <span class="keyword">NVIDIA Nsight Compute</span> instances are active on the same system,
set <samp class="ph codeph">TMPDIR</samp> to a different directory for which the current user has write permissions.
</li>
<li class="li">
Ask the user owning the file, or a system administrator, to remove it or add write permissions for all potential users.
</li>
<li class="li">
On Linux systems setting <samp class="ph codeph">fs.protected_regular=1</samp>,
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 <samp class="ph codeph">sudo sysctl fs.protected_regular=0</samp>, 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="xref" href="https://developer.nvidia.com/ERR_NVGPUCTRPERM" target="_blank" shape="rect">https://developer.nvidia.com/ERR_NVGPUCTRPERM</a> on how to change this setting).
</li>
</ul>
</div>
</li>
<li class="li"><strong class="ph b">Profiling failed because a driver resource was unavailable.</strong><p class="p">
The error indicates that a required CUDA driver resource was unavailable during profiling.
Most commonly, this means that <span class="keyword">NVIDIA Nsight Compute</span> could not reserve the
driver's performance monitor, which is necessary for collecting most metrics.
</p>
<p class="p">
This can happen if another application has a concurrent reservation on this resource.
Such applications can be e.g. <a class="xref" href="https://developer.nvidia.com/dcgm" target="_blank" shape="rect">DCGM</a>,
a client of <a class="xref" href="https://developer.nvidia.com/cupti" target="_blank" shape="rect">CUPTI's Profiling API</a>,
<a class="xref" href="https://developer.nvidia.com/nsight-graphics" target="_blank" shape="rect">Nsight Graphics</a>,
or another instance of <span class="keyword">NVIDIA Nsight Compute</span> without access to the same file system
(see <a class="xref" href="index.html#serialization" shape="rect">serialization</a> for how this is prevented within the same file system).
</p>
<p class="p">
If you expect the problem to be caused by DCGM, consider using <samp class="ph codeph">dcgmi profile --pause</samp> to stop its monitoring
while profiling with <span class="keyword">NVIDIA Nsight Compute</span>.
</p>
</li>
<li class="li">
<p class="p"><strong class="ph b">Could not deploy stock * files to *</strong></p>
<p class="p"><strong class="ph b">Could not determine user home directory for section deployment.</strong></p>
<p class="p">
An error occurred while trying to deploy stock section or rule files.
By default, <span class="keyword">NVIDIA Nsight Compute</span> tries to deploy these to a versioned directory in
the user's home directory (as identified by the <samp class="ph codeph">HOME</samp> environment variable on Linux),
e.g. <samp class="ph codeph">/home/user/Documents/NVIDIA Nsight Compute/<version>/Sections</samp>.
</p>
<p class="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 <span class="keyword">NVIDIA Nsight Compute</span> falls back
to using stock sections and rules from the installation directory.
</p>
<p class="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 <samp class="ph codeph"><span class="keyword">ncu</span> --section-folder-restore</samp>,
or by making <samp class="ph codeph">/home/user/Documents/NVIDIA Nsight Compute/<version></samp>
a symlink to a writable directory.
</p>
</li>
<li class="li">
<p class="p"><strong class="ph b"><samp class="ph codeph">ProxyJump</samp> SSH option is not working</strong></p>
<p class="p"><span class="keyword">NVIDIA Nsight Compute</span> does not manage authentication or interactive prompts with the OpenSSH client launched when using the
<samp class="ph codeph">ProxyJump</samp> 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, <span class="keyword">NVIDIA Nsight Compute</span> should be able to connect to the
target, too.
</p>
</li>
<li class="li">
<p class="p"><strong class="ph b">SSH connection fails without trying to connect</strong></p>
<p class="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 <samp class="ph codeph">IP/Host Name</samp>, <samp class="ph codeph">User Name</samp> and <samp class="ph codeph">Port</samp> fields are correctly set.
</p>
</li>
<li class="li">
<p class="p"><strong class="ph b">SSH connections are still not working</strong></p>
<p class="p">
The problem might come from <span class="keyword">NVIDIA Nsight Compute</span>'s SSH client not finding a suitable host key algorithm to use which
is supported by the remote server. You can force <span class="keyword">NVIDIA Nsight Compute</span> to use a specific set of host key algorithms by
setting the <samp class="ph codeph">HostKeyAlgorithms</samp> 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 <samp class="ph codeph">ssh-keyscan</samp> utility which comes with
the OpenSSH client.
</p>
</li>
<li class="li">
<p class="p"><strong class="ph b">Removing host keys from known hosts files</strong></p>
<p class="p">
When connecting to a target machine, <span class="keyword">NVIDIA Nsight Compute</span> tries to verify the target's host key against the same local
database as the OpenSSH client.
If <span class="keyword">NVIDIA Nsight Compute</span> 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 <samp class="ph codeph">ssh-keygen -R <host></samp> command.
</p>
</li>
</ul>
</div>
</div>
<div class="topic concept nested0" id="notices-header"><a name="notices-header" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#notices-header" name="notices-header" shape="rect">Notices</a></h2>
<div class="topic reference nested1" id="notice"><a name="notice" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#notice" name="notice" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Notice</h3>
<p class="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 class="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>
</div>
</div>
</div>
<div class="topic reference nested1" id="trademarks"><a name="trademarks" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#trademarks" name="trademarks" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Trademarks</h3>
<p class="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>
</div>
</div>
</div>
<div class="topic reference nested1" id="copyright-past-to-present"><a name="copyright-past-to-present" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#copyright-past-to-present" name="copyright-past-to-present" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Copyright</h3>
<p class="p">© <span class="ph">2018</span>-<span class="ph">2022</span> NVIDIA
Corporation and affiliates. All rights reserved.
</p>
<p class="p">This product includes software developed by the Syncro Soft SRL (http://www.sync.ro/).</p>
</div>
</div>
</div>
</div>
<hr id="contents-end"></hr>
</article>
</div>
</div>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/formatting/common.min.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-write.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-tracker.js"></script>
<script type="text/javascript">_satellite.pageBottom();</script>
<script type="text/javascript">var switchTo5x=true;</script><script type="text/javascript" src="http://w.sharethis.com/button/buttons.js"></script><script type="text/javascript">stLight.options({publisher: "998dc202-a267-4d8e-bce9-14debadb8d92", doNotHash: false, doNotCopy: false, hashAddressBar: false});</script></body>
</html>
|