Chcąc pisać w CUDA'ch, tworzymy pliki tekstowe z rozszerzeniem *.cu. Te następnie są on kompilowane z pomocą nvcc (kompilatora dostarczanego przez Nvidię w CUDA Toolkit) do jednego z kilku formatów: (np. PTX, CUBIN, EXE). Pierwszy z nich zawiera instrukcje asemblerowe, które są czytelne dla ludzi (jak ktoś oczywiście lubi / umie). Jest to również uniwersalny format, gdyż przed wykonaniem takiego kodu, jest on kompilowany pod konkretny układ GPU. Natomiast format CUBIN, to już format binarny pod konkretny procesor.
Więcej szczegółów na temat procesu kompilowania i linkowania Cudy zamieszczam na poniższym rysunku (schemat ze strony: http://www.think-techie.com/2009/09/gpu-computing-using-jcuda.html):
W plikach *.cu możemy umieszczać typowy kod C jak i tzw. The Kernel Code. Jest to kod, który wykonuje się na urządzeniu (ang. device) obsługującym CUDA. Kompilator nvcc nie jest przez to samowystarczalny. Do jego poprawnego działania wymagany jest jeszcze kompilator C, np. Visual Studio C++. Początkowo zainstalowałem VC++ 2010 Express Edition, ale się okazało, że moja wersja CUDA współpracuje z MCVC 8.0 lub 9.0. Zainstalowałem więc Visual Studio 2008 i przy próbie kompilacji przykładu Creating kernels dostałem następujący błąd:
nvcc fatal : Cannot find compiler 'cl.exe' in PATH
Ok, szukamy cl.exe I dodajemy do PATH, czyli ***\Microsoft Visual Studio 9.0\VC\bin.
Teraz robi się ciekawie. Uruchamiając przykład, dostaję:
nvcc -m64 -ptx JCudaVectorAddKernel.cu -o JCudaVectorAddKernel.ptx nvcc process exitValue -1 errorMessage: nvcc fatal : Cannot find compiler 'cl.exe' in PATH
ale uruchamiając to samo „z palca” dostaję:
nvcc fatal: Visual Studio configuration file '(null)' could not be found for installation at...
Musiałem doinstalować X64 Compilers and Tools
I zaczęło śmigać.
Czasem może być jeszcze wymagane dodanie do PATH’a: ***\Microsoft Visual Studio 9.0\VC\bin\amd64
Można również się natknąć na poniższy problem:
host_config.h(96) : fatal error C1083: Cannot open include file: 'crtdefs.h': No such file or directory
Aby go rozwiązać, trzeba jeszcze wskazać CUDA, gdzie są pliki nagłówkowe z kompilatora C. Aby to zrobić, należy w pliku nvcc.profile (w katalogu bin instalacji CUDA) w sekcji INCLUDES dodać ścieżkę do odpowiedniego katalogu:
INCLUDES += "-I$(TOP)/include" "-I$(TOP)/include/cudart" "-Id:/Microsoft Visual Studio 9.0/VC/include" $(_SPACE_)
I już działa. Rozwiązanie ostatniego problemu znalazłem tutaj: http://blogs.hoopoe-cloud.com/index.php/2009/09/cudanet-examples-issues
Skoro już wszystko działa (mam nadzieję że niczego nie ominąłem) czas na jakiś kod. Sztandarowym przykładem pokazującym potęgę GPU jest sumowanie wektorów. Poniżej przedstawię przykład ze strony JCuda Code samples. Musiałem jednak trochę go zmodyfikować, gdyż tamtejszy przykład jest przygotowany dla cudów 4.0 i korzysta z metod, których w mojej wersji API (0.3.2a) nie ma. Podrasowałem również kod samego kernela, aby wyjaśnić pewne zjawiska tam zachodzące. Kod biblioteki jest na licencji MIT/X11 License więc mogę modyfikować, z zachowaniem informacji o autorze.
Zacznijmy od kodu Javowego:
// Enable exceptions and omit all subsequent error checks
JCudaDriver.setExceptionsEnabled(true);
// Create the PTX file by calling the NVCC
String ptxFileName = preparePtxFile("JCudaVectorAddKernel.cu");
// Initialize the driver and create a context for the first device.
cuInit(0);
CUdevice device = new CUdevice();
cuDeviceGet(device, 0);
CUcontext context = new CUcontext();
cuCtxCreate(context, 0, device);
// Load the ptx file.
CUmodule module = new CUmodule();
cuModuleLoad(module, ptxFileName);
// Obtain a function pointer to the "add" function.
CUfunction function = new CUfunction();
cuModuleGetFunction(function, module, "add");
Mamy tutaj trochę rzeczy związanych z ustawieniem sterownika Cuda. Początkowo jest włączane rzucanie wyjątków z metod biblioteki, gdy coś pójdzie nie tak. Samo już istnienie takiej metody świadczy według mnie o autorze, któremu jest bliżej do C niż do Javy. Moje obawy zostały potwierdzone przez początkowe formatowanie kodu w źródłowym pliku.
Następnie w linii 5tej ładujemy nasz plik *.cu, a dokładniej, jeżeli nie istnieje dla niego odpowiednik w formacie PTX, to nasz plik źródłowy z kodem kelnera jest kompilowany, w trakcie działania naszej aplikacji. Jest to metoda napisana przez autora przykładu i jak ktoś chce się jej przyjrzeć bliżej to niech zajrzy do kodu.
Wracając do opisanego przykładu dalej następuje inicjalizacja naszego urządzenia i utworzenie dla niego kontekstu. Dalej w linii 15stej nastepuje utworzenie modułu i załadowanie do niego skompilowanego pliku PTX, zwróconego przez metodę preparePtxFile(). W ostatniej linijce w końcu dobieramy się do naszej funkcji, którą napisaliśmy w naszym pliku *.cu.
Dobra teraz czas na przygotowanie danych wejściowych do naszej funkcji:
int numElements = 100000;
// Allocate and fill the host input data
float hostInputA[] = new float[numElements];
float hostInputB[] = new float[numElements];
for (int i = 0; i < numElements; i++) {
hostInputA[i] = (float) i;
hostInputB[i] = (float) i;
}
Tworzymy dwa wektory o długości 100k i wypełniamy kolejnymi liczbami. Zabieg ten ma na celu łatwą weryfikację poprawności uzyskanego wyniku. Stworzone dane wejściowe zostaną utworzone na stercie i nie będą dostępne dla GPU. Po za tym odczyt ten mógłby być bardzo długi (z powodu prędkości magistrali, obecności pamięci wirtualnej i stronicowania). Dla tego względu należy dane przekopiować do obszaru urządzenia. Robotę wykonuje poniższy kod:
// Allocate the device input data, and copy the
// host input data to the device
CUdeviceptr deviceInputA = new CUdeviceptr();
cuMemAlloc(deviceInputA, numElements * Sizeof.FLOAT);
cuMemcpyHtoD(deviceInputA, Pointer.to(hostInputA),
numElements * Sizeof.FLOAT);
CUdeviceptr deviceInputB = new CUdeviceptr();
cuMemAlloc(deviceInputB, numElements * Sizeof.FLOAT);
cuMemcpyHtoD(deviceInputB, Pointer.to(hostInputB),
numElements * Sizeof.FLOAT);
// Allocate device output memory
CUdeviceptr deviceOutput = new CUdeviceptr();
cuMemAlloc(deviceOutput, numElements * Sizeof.FLOAT);
W linii 3ciej tworzymy wskaźnik, który może wskazywać na pamięć urządzenia. Następnie alokujemy do tego wskaźnika pamięć o wielkości numElements * Sizeof.FLOAT. Pierwsza wartość oznacza długość wektora, a druga wielkość zajmowanej pamięci przez jeden element. Sizeof.FLOAT wynosi 4 (bajty).
Następnie w linii 5tej i 6tej następuje kopiowanie ze „zwykłej” pamięci (host) do urządzenia (device), czyli metoda cuMemcpyHtoD(). Najpierw jako argument podajemy wskaźnik do pamięci urządzenia, później referencję dla naszego wektora z pamięci hosta. Na koniec jeszcze ilość bajtów, jaką należy skopiować. Analogiczne operacje są powtarzane dla drugiego argumentu wejściowego. Na koniec jest jeszcze alokowana pamięć na urządzeniu, gdzie będzie zapisywany wynik.
Dobra czas na przekazanie argumentów do funkcji. W nowszym API można to zrobić bardziej przyjemniej. Ja musiałem się zapytać na forum, co z tym fantem, ale na szczęście autor szybko odpisał. Rozwiązanie dla wersji 0.3.2a poniżej:
// sets parametrs
int offset = 0;
offset = align(offset, Sizeof.INT);
cuParamSetv(function, offset, Pointer.to(new int[]{numElements}), Sizeof.INT);
offset += Sizeof.INT;
offset = align(offset, Sizeof.POINTER);
cuParamSetv(function, offset, Pointer.to(deviceInputA), Sizeof.POINTER);
offset += Sizeof.POINTER;
offset = align(offset, Sizeof.POINTER);
cuParamSetv(function, offset, Pointer.to(deviceInputB), Sizeof.POINTER);
offset += Sizeof.POINTER;
offset = align(offset, Sizeof.POINTER);
cuParamSetv(function, offset, Pointer.to(deviceOutput), Sizeof.POINTER);
offset += Sizeof.POINTER;
Jako pierwszy argument przekazujemy do funkcji długość naszego wektora. Jest to liczba typu int i w linii 5tej właśnie ustawiamy, że numElements będzie przekazywane jako pierwszy argument. Jako że nie możemy utworzyć Pointer’a do pojedynczej zmiennej, trzeba zrobić na szybko jednoelementową tablicę i przekazać jako rozmiar wielkość int’a.
Od razu liczony jest offset, który w ostatniej metodzie danego listingu jest przekazywany do funkcji. Pewnie to jest po to, aby mechanizm wywołujący metodę, wiedział ile danych trzeba odczytać ze stosu (moje podejrzenie).
Dobra, mamy już dane wejściowe, mamy je przyporządkowane jako argument funkcji, najwyższy czas na wywołanie:
// Call the kernel function. int blockSizeX = 256; int gridSizeX = (int) Math.ceil((double) numElements / blockSizeX); cuParamSetSize(function, offset); // run grid cuLaunchGrid(function, gridSizeX, blockSizeX); cuCtxSynchronize();
Ustalamy wielkość bloku i gridu, który posłuży nam do obliczeń, przekazujemy długość parametrów i wywołujemy funkcję. O co chodzi z blokami i gridami to możecie doczytać w książce CUDA w przykładach lub w CUDA C Programming Guide (obrazek 2-1, strona 9) który jest również dostępny po zainstalowaniu CUDA Toolkit.
Metodę cuCtxSynchronize() wywołujemy po to, aby w razie niepowodzenia jak najszybciej się dowiedzieć, gdzie mniej więcej jest błąd. Pozostaje jeszcze skopiowanie wyniku:
// Allocate host output memory and copy the device output
// to the host.
float hostOutput[] = new float[numElements];
cuMemcpyDtoH(Pointer.to(hostOutput), deviceOutput,
numElements * Sizeof.FLOAT);
do RAMu, porównanie wyników:
// Verify the result
boolean passed = true;
for (int i = 0; i < numElements; i++) {
float expected = i + i;
if (Math.abs(hostOutput[i] - expected) > 1e-5) {
System.out.println(
"At index " + i + " found " + hostOutput[i] +
" but expected " + expected);
passed = false;
break;
}
}
System.out.println("Test " + (passed ? "PASSED" : "FAILED"));
i sprzątanie:
// Clean up. cuMemFree(deviceInputA); cuMemFree(deviceInputB); cuMemFree(deviceOutput);
No to czas na kod wykonywany na GPU:
extern "C"
__global__ void add(int n, float *a, float *b, float *sum)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
while(i < n) {
sum[i] = a[i] + b[i];
i = i + blockDim.x * gridDim.x;
}
}
Linijkę extern "C", potrzebujemy, aby kompilować naszą funkcję zgodnie z regułami kompilowania dla ANSI C. Następnie __global__ powoduje, że metoda oznaczona takim kwalifikatorem, będzie się wykonywać na GPU i może być wywołana z poziomu hosta (CPU). Inna użyteczna definicja to __device__. Funkcje oznaczone w ten sposób, również działają na GPU, ale mogą jedynie być wywołane z poziomu innych funkcji GPU.
Dalej mamy sygnaturę metody, czyli przekazujemy kolejno długość wektorów i wskaźniki na wektory wejściowe i jeden na wyjściowy. Kod jest podobny do tego przedstawionego w CUDA w przykładach (str. 66). W pierwszej linijce funkcji wyliczamy indeks, od którego będziemy zaczynać sumowanie. blockDim reprezentuje liczbę wątków w każdym wymiarze bloku. My stosujemy tylko jednowymiarowy blok. W blockIdx znajduje się indeks bloku, który aktualnie wykonuje kod danej funkcji, a w threadIdx indeks aktualnego wątku. Są to 3 standardowe zmienne systemu wykonawczego CUDA. Siatki bloków mogą być dwuwymiarowe, a każdy blok może mieć trójwymiarową tablicę wątków. Więc jak się ktoś orientuje w pięciu wymiarach, to ma pole do popisu.
Dalsze obliczenia (pętla while) wykonujemy od póki indeks jest mniejszy od długości zadanego wektora. W linijce 6tej nic się ciekawego nie dzieje (po prostu jest obliczany wynik) i w następnie obliczamy kolejny indeks. Normalnie (w przypadku implementacji jednowątkowej) to byśmy wykonywali inkrementację, ale jako że kod jest wykonywany na GPU, to kolejny indeks musimy trochę inaczej policzyć. Zwiększamy go o iloczyn liczby wątków na blok i liczby bloków w siatce.
Na razie to tyle co przygotowałem o cudach. Kod znajdziecie na githubie (JCudaExample), choć nie jest on nie wiadomo jak odkrywczy. Przy uruchomieniu mogą pojawić się kłopoty, więc warto dopasować ścieżkę (java.library.path) w konfiguracji uruchomieniowej.
Podsumowując, jeśli mamy sporo równoległych obliczeń do wykonania, dobrą kartę graficzną i zależy nam na szybkim czasie działania, to można się pobawić w Cuda. Nie jest to jednak tak przyjemnie jak JVM, gdzie wiadomo, z której linijki leci wyjątek i wszystko staje się szybciej jasne. Mimo wszystko praca ze sprzętem daje dużo satysfakcji.
