NVIDIA Nsight - podstawy

Wprowadzenie¶

Niniejszy tutorial pokazuje podstawy korzystania z narzędzi Nsight Systems oraz Nsight Compute w celu analizy wydajności programów CUDA. Zarys możliwości wspomnianych narzędzi jest przedstawiony na stronie NVIDIA Nsight.

Będziemy pracować na przykładzie prostego programu CUDA którego zadaniem jest konwolucja obrazu. Jest to technika przetwarzania obrazów polegająca na nakładaniu filtrów w wyniku których powstaje nowy obraz, a wartość każdego z pikseli jest kombinacją liniową jego pierwotnych sąsiadów. Metoda ta jest stosowana na przykład do wykrywania krawędzi lub dodawania rozmycia.

Skupimy się na podstawowych metrykach i aspektach aplikacji jakie można zbadać za pomocą narzędzi Nsight. Postaramy się jednak wskazać miejsca, w których istnieje możliwość pogłębionej analizy.

Wymagania¶

Aby wykonać poniższy tutorial niezbędny jest dostęp do komputera:

wyposażonego w kartę GPU NVIDIA
posiadającego zainstalowany CUDA Toolkit (samodzielny lub jako element NVIDIA HPC Toolkit, informacje o pakietach - patrz CUDA > Informacje o standardzie) w wersji 11.0 lub wyższej oraz wspierającej używaną kartę GPU

Obydwa powyższe toolkity zawierają środowisko i kompilator niezbędne do wykonywania programów CUDA oraz narzędzia Nsight Systems i Nsight Compute.

Do kompilacji programu wymagany jest:

dostęp do biblioteki OpenCV,
system budowania CMake.

Jeżeli na urządzeniu z kartą GPU nie mamy możliwosci uruchamiania aplikacji okienkowych, analizy zebranych profili aplikacji można dokonać na innej maszynie. W takim przypadku:

na komputerze wyposażonym w GPU wykonujemy program (potrzebujemy tam wyżej wymienione środowisko),
na drugą maszynę pobieramy profile i tam otwieramy je w trybie graficznym w narzędziach Nsight.

Oczywiście drugi komputer musi posiadać zainstalowane narzędzia Nsight Systems oraz Compute (w ramach pakietu CUDA Toolkit lub indywidualnie, patrz NVIDIA Nsight > dostępność).

Przykładowy program¶

Kod przykładowego programu wykonującego konwolucję pojedynczego obrazu, wraz z przykładowym obrazem wejściowym (o rozdzielczości 6048x4024) i wynikowym można pobrać w formie archiwum cuda-image-convolution.zip o następującej zawartości:

simple-convolution.cu
CMakeLists.txt
source.jpg
example-out.jpg

Kod źródłowy – funkcja main

int main() {
    const uint16_t block_size = 1;
    const uint8_t filter_size = 3;
    const std::vector<float> filter{-1, -1, -1,
                                    -1, 8, -1,
                                    -1, -1, -1};

    run_simple("source.jpg", cv::IMREAD_GRAYSCALE, filter.data(), filter_size, block_size);

    return 0;
}

W funkcji main zdefiniowane są trzy stałe programu: rozmiar bloku w gridzie obliczeniowym, a także rozmiar oraz zawartość macierzy filtra konwolucji (tzw. kernel). Są one następnie, wraz z nazwą obrazu wejściowego oraz trybem odczytu (patrz dokumentacja OpenCV), przekazywane jako argumenty funkcji run_simple. Funkcja ta jest odpowiedzialna za przeprowadzenie obliczeń: wczytuje obraz, alokuje pamięć GPU, dokonuje transferu danych, po czym wywołuje właściwy kernel, który wykona obliczenia na GPU.

Kod źródłowy – kernel (konwolucja)

__global__ void simple_conv(const uint8_t *img_in, uint8_t *img_out,
                        uint32_t img_width, uint32_t img_height, uint8_t img_channels,
                        const float *filter, uint8_t filter_size) {

    const int col = threadIdx.x + blockIdx.x * blockDim.x;
    const int row = threadIdx.y + blockIdx.y * blockDim.y;

    if (row < img_height and col < img_width) {
        for (auto i_channel = 0; i_channel < img_channels; ++i_channel) {
            const int row_start = row - 1;
            const int col_start = col - 1;

            float accum = 0;
            for (auto i_row = 0; i_row < filter_size; ++i_row) {
                for (auto i_col = 0; i_col < filter_size; ++i_col) {

                    const int row_curr = row_start + i_row;
                    const int col_curr = col_start + i_col;

                    if (row_curr >= 0 and row_curr < img_height and col_curr >= 0 and col_curr < img_width) {
                        accum += img_in[(row_curr * img_width + col_curr) * img_channels + i_channel] *
                                 filter[i_row * filter_size + i_col];
                    } else {
                        accum = 0;
                    }
                }
            }
            img_out[(row * img_width + col) * img_channels + i_channel] = min(static_cast<int>(accum), 255);
        }
    }
}

Do zaimplementowania konwolucji wykorzystaliśmy model programowania CUDA. Każdy wątek (thread) wykonuje obliczenia dla pojedynczego piksela: nowa wartość obliczana jest jako suma elementów macierzy filtra pomnożonych przez wartości odpowiednich sąsiadów lub wartość początkową.

Do kompilacji programu korzystamy z systemu budowania CMake. W podstawowej konfiguracji CMake w systemie Linux, aby skompilować aplikację trzeba wejść do katalogu z rozpakowanym archiwum i wywołać dwie komendy:

cmake .
make

W wyniku otrzymujemy wykonywalny plik binarny simple-convolution, który można uruchomić bez dodatkowych parametrów:

./simple-convolution

Program ten przetworzy załączony obraz source.jpg i wykona na nim konwolucję z filtrem wykrywania krawędzi zadanym poniższą macierzą:

$\begin{bmatrix}-1 & -1 & -1\\ -1 & 8 & -1 \\ -1 & -1 & -1\end{bmatrix}$

W efekcie wykonania programu w katalogu roboczym zostanie zapisany plik wynikowy out_source.jpg z przetworzonym obrazem.

obraz wyjściowy po zastosowaniu filtra wykrywania krawędzi

Tutorial¶

W środowiskach HPC, dostęp do trybu graficznego nie zawsze jest możliwy lub jest limitowany czasowo. Oprócz tego, nawet gdybyśmy mieli dostęp do GUI to analiza zwykle wymaga sporej ilości naszego czasu pracy, w którym to zasoby obliczeniowe nie są wykorzystywane. Z tego powodu skupimy się na uruchamianiu narzędzi Nsight z poziomu wiersza polecenia i późniejszym importowaniu wyników do wersji graficznej.

Jak wspominaliśmy we wprowadzeniu, taki schemat pracy pozwala na ewentualne prowadzenie analizy na innej maszynie z dostępnym oprogramowaniem Nsight. W przypadku systemów HPC, zbieranie profili można wtedy wykonywać w ramach zadań wsadowych, które nie wymagają interakcji z użytkownikiem.

Komendy wiersza poleceń

W pojawiających się dalej komendach wiersza poleceń zakładamy, że ścieżki do plików wykonywalnych Nsight Systems oraz Nsight Compute znajdują sią w zmiennej środowiskowej PATH, czyli że programy te można uruchamiać bezpośrednio wpisując w konsoli polecenia nsys lub ncu.

Dodatkowo, wszystkie komendy są wywoływane z katalogu zawierającego kod źródłowy przykładowego programu.

Testowa konfiguracja

Wszystkie prezentowane poniżej operacje zostały wykonane na maszynie:

działającej pod kontrolą systemu Linux,
wyposażonej w GPU NVIDIA RTX 4090,
z CUDA Toolkit w wersji 11.8,
oraz narzędziami Nsight w wersji 2023.1.

Krok 1: Nsight Systems – podstawowa analiza¶

Na początek spróbujmy uzyskać ogólny obraz zachowania programu przeprowadzającego konwolucję. Wykorzystamy do tego narzędzie Nsight Systems.

Uruchamiamy profilowanie programu za pomocą Nsight Systems

nsys profile ./simple-convolution

W efekcie tej komendy nastąpi wykonanie programu oraz otrzymamy raport o domyślnej nazwie report1.nsys-rep (w starszych wersjach Nsight Systems, przed 2021.4, będzie to report1.qdrep). Domyślną konfigurację profilowania można zmienić dodatkowymi opcjami (zobacz nsys profile --help).

Output w konsoli oraz możliwe błędy

W wyniku wykonania komendy nsys profile w konsoli zobaczymy mniej więcej następujące komunikaty:

Collecting data...
Processing events...
Capturing symbol files...
Saving temporary "<...>/tmp/nsys-report-<...>.qdstrm" file to disk...
Creating final output files...
Processing [============================100%]
Saved report file to "<...>/tmp/nsys-report-<...>.qdrep"
Report file moved to "<...>/simple-convolution/report1.qdrep"

⚠️ W przypadku niekompatybilności z obecną w systemie wersją biblioteki glibc (dotyczy głównie wersji Nsight Systems starszych niż 2021.5 uruchamianych na nowszych wersjach dystrybucji Linuxa) możliwe jest otrzymanie raportu, który nie będzie zawierał informacji o wykonaniu programu na karcie GPU. Możemy rozpoznać taką sytuację wtedy gdy otrzymamy w konsoli np. poniższy komunikat:

./simple-convolution: symbol lookup error: <...>/nsight-systems-2021.1.3/target-linux-x64/libToolsInjectionProxy64.so: undefined symbol: __libc_dlsym, version GLIBC_PRIVATE

W takiej sytuacji należy sięgnąć po nowszą wersję Nsight Systems.

Włączamy graficzną wersję Nsight Systems

nsys-ui

Uruchamianie profilowania z poziomu GUI

Alternatywnie, profilowanie może być także uruchomione bezpośrednio z poziomu okna aplikacji – zarówno lokalnie, jak i zdalnie poprzez połączenie SSH. Należy w tym celu utworzyć nowy projekt (File -> New Project) oraz wybrać target, czyli komputer lokalny lub połączenie zdalne, na którym ma zostać uruchomiona profilowana aplikacja. W kolejnym oknie należy wskazać plik wykonywalny, katalog roboczy (zawierający plik wykonywalny i/lub dane wejściowe) oraz, ewentualnie, dostosować opcje profilowania.

nsys-ui — domyślna konfiguracja profilowania po wybraniu *targetu*

Otwieramy raport (File -> Open)

Bezpośrednio po otworzeniu powinien on wyglądać podobnie do przedstawionego poniżej. Jest to widok osi czasu (Timeline View). Raport można również otworzyć bezpośrednio podając nsys-ui report1.nsys-rep.

Poruszanie się po interfejsie Nsight Systems może ułatwić lista skrótów klawiszowych (1). Informacje o wykonaniu aplikacji oraz przebiegu profilowania mogą pojawić się na liście wiadomości (2). Kliknięcie w to miejsce przenosi z widoku osi czasu (Timeline View) na widok diagnostyki (Diagnostics Summary), który może być pomocny w zdiagnozowaniu ewentualnych problemów z wykonaniem profilowania. W naszym wypadku nie zawiera on jednak kluczowych komunikatów, więc go pomijamy.

Widok osi czasu przedstawia poszczególne etapy działania aplikacji zarówno po stronie hosta (CPU) jak i GPU. W górnej części znajduje się oś reprezentująca czas działania (3). Poniżej widzimy różne zdarzenia. Najważniejsze dla nas są elementy bezpośrednio związane z API CUDA: alokacja pamięci (4) oraz wykonanie kernela i transfery pamięci (5).

Już na tym etapie widać, że w przypadku niewielkich obliczeń, jak przeprowadzenie konwolucji na pojedynczym obrazie, wykonanie właściwych obliczeń (kernela) to niewielki ułamek czasu działania aplikacji. Stosunkowo dużo czasu pochłania w takim wypadku sama inicjalizacja środowiska CUDA i operacje takie jak alokacja i transfery pamięci.

Przybliżamy (Ctrl + scroll) widok kernela (5)

Rozwijamy szczegóły sekcji "CUDA HW (...)" (1)

Skupmy się jednak na samym kernelu, a więc operacji konwolucji. Wykonanie kernela wraz z powiązanymi transferami pamięci (memcpy) jest teraz dokładnie widoczne na osi czasu.

Zmieniamy dolny widok z "Events View" na "Stats System View" oraz wybieramy zakładkę "CUDA GPU Kernel/Grid/Block Summary" (2)

W dolnej sekcji raportu w widoku statystyk widzimy teraz informacje o parametrach wykonania kernela simple_conv (3).

W widoku "Stats System View" wybieramy zakładkę "CUDA GPU Summary (Kernel/MemOps)" (1)

W widoku statystyk można znaleźć także inne istotne informacje takie jak precyzyjne dane o czasie trwania wykonania kerneli oraz transferów pamięci (2). Przypomnijmy, że memcpy DtoH oznacza kopiowanie z pamięci GPU do pamięci RAM (device-to-host), a HtoD oznacza kopiowanie z RAM do GPU (host-to-device).

Warto pamiętać, że szczegóły raportu, takie jak dokładne czasy trwania poszczególnych etapów, mogę różnić się w zależności od użytej karty GPU.

Wprawne oko już na podstawie powyższych statystyk może wskazać pierwsze elementy, które można zoptymalizować. Aby jednak uzyskać dokładniejszy obraz wykonania aplikacji, zobaczymy jaki będzie wynik analizy za pomocą narzędzia Nsight Compute.

Krok 2: Nsight Compute - szczegółowa analiza¶

Uprawnienia wymagane do profilowania z użyciem Nsight Compute

Nsight Compute wymaga aby użytkownik który uruchamia profilowanie miał dostęp do liczników zdarzeń na karcie GPU (tzw. GPU performance counters). Jeśli ich nie ma, może zobaczyć błąd ERR_NVGPUCTRPERM. W praktyce, konfiguracja dostępu do liczników wymaga uprawnień administratora systemu. Więcej szczegółów ERR_NVGPUCTRPERM: Permission issue with Performance Counters.

Pierwsze uruchomienie¶

Do bardziej szczegółowej analizy, a także do pomocy przy optymalizacji omawianego przez nas programu wykorzystamy narzędzie Nsight Compute.

Uruchamiamy profilowanie programu za pomocą Nsight Compute

ncu -o profile ./simple-convolution

Uruchamianie profilowania z poziomu GUI

Podobnie jak w przypadku Nsight Systems, możliwe jest uruchomienie profilowania z poziomu interfejsu graficznego. Aby rozpocząć profilowanie należy utworzyć nowy projekt (File -> New Project), nadać mu nazwę oraz wybrać lokalizację. Następnie wybieramy target oraz pozostałe opcje w sposób analogiczny z Nsight Systems.

Jako wynik działania profilera otrzymujemy raport o domyślnej nazwie profile.ncu-rep.

Uruchamiamy wersję okienkową Nsight Compute

ncu-ui

Otwieramy wygenerowany raport (File -> Open File lub bezpośrednio za pomocą komendy ncu-ui profile.ncu-rep).

W naszym wypadku raport po otworzeniu powinien przedstawiać się podobnie do poniższego.

Domyślnym (startowym) widokiem jest widok szczegółów kerneli (1). W naszym programie wykonuje się tylko jeden kernel, jest on więc domyślnie wybrany do podsumowania (2), a pod spodem widoczne są szczegóły jego wykonania podzielone na kilka sekcji. Przykładowo, w pierwszej części - GPU Speed of Light Throughput - możemy znaleźć informacje takie jak czas wykonania kernela (3) czy stopień wykorzystania GPU oraz poszczególnych rodzajów pamięci (4).

Pod każdą z sekcji podsumowania otrzymaliśmy po jednej sugestii poprawek pod kątem wydajności. Stwierdzenie, które wskazówki mogą przynieść najlepsze efekty nie zawsze jest oczywiste, natomiast korekty dotyczące launch statistics (5) są często stosunkowo proste do wprowadzenia. W naszym wypadku uwaga dotyczy zbyt małego block size (6), co skutkuje bardzo niskim wykorzystaniem procesora GPU. Nsight Compute sugeruje takie skonfigurowanie gridu obliczeniowego, aby rozmiar bloku był wielokrotnością liczby 32.

Optymalizacja rozmiaru bloku¶

Zobaczmy więc, jaki będzie rezultat uruchomienia tego samego kernela, ale skonfigurowanego z rozmiarem bloku wynoszącym 32, co przełoży się na bloki CUDA o wymiarach 32x32.

Zmieniamy wartość stałej block_size w funkcji main na 32

Po ponownym skompilowaniu aplikacji i uruchomieniu profilowania otrzymujemy następujący rezultat.

Jak widać, dzięki lepszej konfiguracji zrównoleglenia, czas wykonania kernela skrócił się z ponad 13 ms do niecałych 400 us (1), a więc ponad 30-krotnie!

Poprawa zajętości warpów¶

Pozostaje jeszcze jednak istotna uwaga o zajętości warpów (2), która w przypadku obecnej konfiguracji mogłaby maksymalnie wynieść 67% (Theoretical Occupancy) a w praktyce wyniosła 46% (Achieved Occupancy). Zasady doboru parametrów w celu uzyskania jak najlepszej zajętości zależą od compute capability używanego GPU oraz samego kernela (wykorzystywanych rejestrów) i są szczegółowo opisane w dobrych praktykach CUDA. Spróbujmy jeszcze raz zmienić block size tak aby zmniejszyć liczbę rejestrów wykorzystywanych przez pojedynczy blok i zwiększyć liczbę bloków.

Ponownie zmieniamy wartość stałej block_size, zmniejszamy ją do 16

Po kolejnym uruchomieniu profilowania widać poprawę zajętości teoretycznej i osiągniętej do odpowiednio 100% i 76% (1), a także dalsze przyspieszenie kernela konwolucji (2). Pozostałe wygenerowane przez Nsight Compute komentarze odnoszą się już tylko do bardziej zaawansowanych metod optymalizacji.

Podsumowanie¶

W tym tutorialu pokazaliśmy jak poprzez prostą analizę można rozpocząć proces optymalizacji kodu akcelerowanego na GPU. Dzięki narzędziom Nsight Systems oraz Nsight Compute udało się przyjrzeć poszczególnym etapom programu wykonującego konwolucję obrazu, a także dokonać prostych, ale skutecznych optymalizacji.

W ten sposób, na bazie analizy bardzo prostego programu, zarysowaliśmy możliwości narzędzi z pakietu Nsight. Ich funkcjonalności jednak na tym się nie kończą, a profilowanie bardziej złożonych aplikacji może wymagać zastosowania dodatkowych technik. Więcej szczegółów na temat profilerów Nsight Systems oraz Nsight Compute można znaleźć na przykład w następujących materiałach wideo:

Ostatnia aktualizacja: 1 grudnia 2023