Dziś troszkę z innej beczki. Mianowicie będzie trochę o wydajności i równoległym przetwarzaniu danych. No to do dzieła!

Karta graficzna to jeden z najistotniejszych komponentów w komputerze. Odpowiada nie tylko za generowanie grafiki 3D, ale także do wspomagania albo wyręczania procesora z pewnych obliczeń. W ostatnich latach na rynku pojawiło się wiele nowych rozwiązań układów graficznych, jak GMA. Intela czy nowe generacje GPU przygotowane przez firmy NVIDIA i AMD. Nowoczesne karty graficzne to wielordzeniowe, wyspecjalizowane procesory, które mogą służyć nie tylko do obliczeń związanych z grafiką. Nowe możliwości GPU dostrzegł koncern NVIDIA, tworząc API rozszerzającego możliwości języka C. Ta technologia pozwala zaprojektować programiście aplikację do rozwiązania bardzo złożonych obliczeniowo zadań w krótkim czasie poprzez wykorzystanie mocy przetwarzania równoległego wielu rdzeni procesora graficznego. Ważny jest również nieograniczony dostęp do przestrzeni adresowej karty i możliwość skorzystania ze współdzielonej pamięci. Technologia ma pełne wsparcie dla operacji logicznych i wykonywalnych na liczbach całkowitych, wektorach i macierzach.

Jednym z problemów, które można rozwiązać za pomocą CUDA (pełna lista na końcu wpisu) jest problem wydajnego mnożenia macierz rzadkiej przez wektor. Macierz to określony zbiór uporządkowanych wierszy i kolumn.  Wektor natomiast może być – w zależności od interpretacji – określony jako macierz jednowierszowa bądź jednokolumnowa. Wynikiem mnożenia macierzy przez wektor jest wektor. A teraz co to jest macierz rzadka. To specyficzna macierz, której najczęściej występującymi.

Dlaczego wykonywanie takiej operacji przez jest czasochłonne? Otóż procesor to jednostka pracująca sekwencyjnie, co oznacza, że wykonuje jedną operację w danej chwili (oczywiście mowa tu o procesorze z jednym rdzeniem). Można powiedzieć, że CUDA to taki procesor o dużej liczbie słabych procesorów, dzięki czemu można w nich wykonywać wiele takich samych operacji atomowych (niepodzielnych) jednocześnie . Fragmenty kodu do wykonania na karcie graficznej muszą być pogrupowane w funkcje (zwane Kernelami) i skompilowane za pomocą dostarczanego przez nVidię kompilatora nvcc lub korzystając z odpowiednich bibliotek do MS Visual Studio. Dozwolone jest także wywoływanie kodu z poziomu C++ czy C#.

Z pośród wielu sposobów prezentacji macierzy rzadkich, zaprezentuję format CSR(Compressed Sparse Row). Charakteryzuje się on zapisem danych w postaci trzech tabel: data, indices i ptr. W pierwszej z nich (data) umieszczone są wszystkie niezerowe wartości w kolejności od lewej strony do prawej, wiersz po wierszu (rozpoczynając od góry). Następna tablica (indices) zawiera numery kolumn niezerowych elementów macierzy, w takim samym porządku jak poprzednia. Ostatni zbiór (ptr) określa ilość cyfr, które nie są zerem w podanym wierszu (pierwsza wartość to zawsze liczbami są zera).

Teraz jak wygląda sama praca karty graficznej. Procesor (Host)zleca wykonanie funkcji (Kernel) na karcie graficznej (Device). Każde wywołanie Kernela przez kod wykonujący się na CPU tworzy na karcie graficznej nowy Grid wątków. Wątki w Gridzie są pogrupowane w bloki o identycznym rozmiarze i indeksowane poprzez pozycję w bloku i pozycję bloku w Gridzie. Wymiary bloku i wymiary Gridu podaje się jako argumenty wywołania Kernela. Możliwe jest wywołanie kilku Kerneli jednocześnie (różnych bądź takich samych) o różnych parametrach wywołania, należy jedynie uważać aby nie zapisywały danych w to samo miejsce w pamięci. Wątki można synchronizować tylko w obrębie bloku. Oznacza to, że nie możliwa jest bezpieczna komunikacja pomiędzy wątkami należącymi do dwóch różnych bloków. Wątki z każdego bloku są dzielone na grupy o stałym rozmiarze równym nazywane Warp. W danej chwili każdy multiprocesor wykonuje jeden warp wątków należących do jednego bloku.

To teraz powrócę do naszego problemu mnożenia macierzy przez wektor. Programowo wygląda to tak:

__global__ void addKernel(CMatrixCSR mat)
{
    int row = blockDim.x * blockIdx.x + threadIdx.x;
                if (row < mat.n ){
                                float dot = 0;
                                int row_start = mat.ptr[row];
                                int row_end = mat.ptr[row + 1];
                                for (int jj = row_start; jj < row_end; jj++)
dot += mat.data[jj] * mat.vector[mat.indices[jj]];

// jeden blok wykonuje 512 sekwencyjnych obliczen
                                mat.solution[row] += dot;
                }
}

Program wydaje się zrozumiały. Kernel przetwarza kolejne wiersze znajdujące się w macierzy. Dla każdego z nich wykonuje standardową operację mnożenia i sumowania.  Należy pamiętać o zaprogramowaniu odpowiedniej struktury przechowującej dane oraz metody odpowiedzialne za wczytywanie i zapisywanie macierzy do pliku. Niestety takie wykonywanie obliczeń jest bardzo czasochłonne i mało wydajne. Spójrzmy więc na to rozwiązanie:

 __global__ void addKernel (CMatrixCSR mat)
{
 __shared__ float vals[512];    // zadeklarowanie tablicy w pamięci współdzielonej dla kazdego bloku w watku. 
512 watkow na blok * (4 bajty/element) = 2 KB    

    int thread_id = blockDim.x * blockIdx.x + threadIdx.x;     //   nadaje unikalny id kazdemu watkowi
    int warp_id = thread_id / 32; //   indeks 'warpu' globalnego
    int lane = thread_id  &(32 - 1); //   indeks wątku w warpie thread   
    
    //   jeden warp przypada na  jeden wiersz (row)
    int row = warp_id;
    
    if  (row < mat.n)
    {
        int row_start = mat.ptr[row];
        int row_end = mat.ptr[row + 1];
        
        //   wykonuje obliczenia sumy
        vals [threadIdx.x] = 0; // czyszczenie bufora
        for (int jj = row_start + lane; jj < row_end; jj += 32)
            vals[threadIdx.x] += mat.data[jj] * mat.vector[mat.indices[jj]];
        
        //  redukacja w pamiêci współdzielonej parallel (współbieżność)
        if (lane < 16) vals[threadIdx.x] += vals[threadIdx.x + 16];    // (16 obliczen wykonywanych wspolbieznie)
        __syncthreads();
        if (lane < 8) vals[threadIdx.x] += vals[threadIdx.x + 8];
        __syncthreads();
        if (lane < 4) vals[threadIdx.x] += vals[threadIdx.x + 4];
        __syncthreads();
        if (lane < 2) vals[threadIdx.x] += vals[threadIdx.x + 2];
        __syncthreads();
        if (lane < 1) vals[threadIdx.x] += vals[threadIdx.x + 1];
        __syncthreads();

        //   pierwszy wątek wypisuje rezultat
        if  (lane == 0)
        mat.solution[row] += vals[threadIdx.x];
    }
}

W ten o to sposób  poprawiliśmy wydajność algorytmu i przyśpieszyliśmy obliczenia wykonujące mnożenie macierzy rzadkich przez wektor. Należy pamiętać o ograniczeniach pamięci przy wczytywaniu danych z pliku (będzie to ok. 4 GB).

 

Inne przykłady wykorzystania CUDA:

• Przyspieszenie szyfrowania i kompresji oraz konwersji wideo do różnych formatów.

• Symulacje fizyczne (np. w dynamice płynów) i obliczenia inżynierskie.

• Obrazowanie wirtualnej rzeczywistości na podstawie tomografii komputerowej i rezonansu magnetycznego.

• Efekty specjalne w grafice komputerowej, np. symulacje falujących powierzchni ubrań.

• Sztuczna inteligencja.

 

blog comments powered by Disqus