wtorek, 27 października 2009

CUDA ruszyła

Pierwszy projekt z wykorzystaniem CUDA mam już za sobą. Sam kod nie był trudny do napisania. Najgorsze było skłonienie Visual Studio do podpowiadania kodu (co do tej pory mi się nie udało - podpowiada kod tylko dla zmiennych i funkcji nieCUDA). Kolejna sprawa, która mnie irytowała to ścieżki dla plików nagłówkowych i bibliotek. Na początku zakładałem projekt w Moje Dokumenty\Visual Studio 2008\Projects i usiłowałem w ustawieniach projektu dodać ścieżki do CUDA SDK. Kompilator wywalał błędy, że nie może znaleźć pliku cuda.h podczas gdy wydawało mi się, że wszystko jest w porządku. Ostatecznie wyszło na to, że następujące ustawienia są prawidłowe (dla konfiguracji Debug):
- Linker -> General -> Additional Library Directories: "C:\CUDA\lib";"C:\Documents and Settings\Yelonek\Ustawienia lokalne\dane aplikacji\NVIDIA Corporation\NVIDIA CUDA SDK\common\lib" (tutaj taka dziwna ścieżka, bo wybrałem "instaluj tylko dla mnie" podczas instalacji pakietu i nie chciało mi się tego odkręcać)
- Linker -> Input -> Additional Dependencies: cudart.lib cutil32d.lib
Następny problem jaki się pojawił to to, że VS nie koloruje składni w plikach cu. Można to obejść w następujący sposób:
W katalogu NVIDIA CUDA SDK\doc\syntax_highlighting znajduje się plik usertype.dat, który należy umieścić w katalogu Microsoft Visual Studio 8\Common7\IDE. Następnie w samym VS należy w Tools... -> Options otworzyć Text Editor w liście po lewej i wybrać File Extension. Wpisujemy cu w polu Extension, ustawiamy edytor na MS Visual C++ i klikamy Add. Po zatwierdzeniu zmian i restarcie programu można się już cieszyć pokolorowaną składnią przy przeglądaniu plików cu.
Istnieje jeszcze jeden sposób na trochę wygodniejsze pisanie programów. Patent ten można znaleźć na innych stronach i blogach, ale też go tu opiszę. Polega on na tym, że program piszemy w pliku C, a podczas budowania projektu plik c kopiuje się do pliku cu, a dopiero plik cu jest budowany za pomocą nvcc. Dokonać tego można definiując własne zasady budowania plików. Po kliknięciu ppm na projekt w Solution Explorerze wybieramy Custom Build Steps... Tworzymy nowy plik zasad (New Rule File), podajemy nazwę jaka będzie wyświetlana w VS, nazwę pliku oraz ściężkę w jakiej chcemy go zapisać. Dodajemy nową zasadę (Add Build Rule). Pierwsza zasada to kopiowanie pliku c do cu.
Command line:copy $(InputFileName) $(InputName).cu
Display Name: Zamiana pliku c na cu
Execution Description: $(InputFileName) ------> $(InputName).cu
File Extensions: *.c
Name: Zamiana pliku c na cu
Outputs: $(InputName).cu
Show Only Rule Properties: True
Supports File Batching: False
Druga zasada to kompilacja plików cu za pomocą nvcc:
Additional dependencies: $(CUDA_INC_PATH)";"../../common/inc"
Command Line: "$(CUDA_BIN_PATH)\nvcc.exe" -ccbin "$(VCInstallDir)bin" -c -D_DEBUG -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/Od,/Zi,/RTC1,/MTd -I"$(CUDA_INC_PATH)" -I../../common/inc -o $(ConfigurationName)\$(InputName).obj $(InputFileName)
Display Name: Kompilacja NVIDIA C for CUDA
Execution Description: Wywoływanie NVCC
File Extension: *.cu
Name: Kompilacja plików CUDA
Outputs: $(ConfigurationName)\$(InputName).obj
Show Only Rule Properties: True
Supports File Batching: False
Teraz tworzymy plik main.c w którym piszemy kod aplikacji. Kopiujemy go (tylko za pierwszym razem) i nazywamy main.cu, a następnie dodajemy do projektu. We właściwościach pliku main.c zamieniamy Tool z C/C++ Compiler na "Zamiana c na cu", a we właściwościach plik main.cu wybieramy "Kompilacja NVIDIA C for CUDA". Od tej pory można już budować projekt za pomocą F6 bez żadnych problemów.
Należy przy tym pamiętać, aby zmian dokonywać w pliku C, a nie CU!!! Gdyby kompilacja projektu się nie udała, to dostaniemy błędy tak jak zawsze, ale dwukrotne kliknięcie przeniesie nas do pliku CU, a nie C!!! Jeśli o tym zapomnimy i zmienimy plik cu, a następnie ponownie zbudujemy projekt wszystko będzie działać: plik c jest niezmieniony, więc VS go ponownie nie buduje (więc plik c nie zostanie skopiowany na plik cu). Załóżmy teraz, że zamknęliśmy projekt, otwieramy go ponownie i edytujemy plik c. Co się dzieje? Plik cu został nadpisany, znika kod pracowicie przez nas napisany i ponownie pojawiają się błędy kompilatora z którymi walczyliśmy przed zamknięciem programu. Można tak stracić parę godzin pracy, więc radzę uważać.

Na koniec przykładowy kod:
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>

// Struktura danych pojedynczego agenta
struct Agent
{
float x;
float y;
};

// Kernel który będzie wykonywał się na karcie
__global__ void Fitness(Agent* agent, int N, float* results)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N)
{
results[idx] = 2 * agent[idx].x * agent[idx].x - agent[idx].x + 3;
}
}

// Funkcja ktora bedzie wykonywac sie na hoscie
int main(void)
{
// Deklaracja zmiennych
const int N = 10000; // Ilosc agentow
Agent *host_agents = NULL, *device_agents = NULL; // Wskaznik do tablic u hosta (_h) i na karcie (_d)
size_t size_agents = N * sizeof(Agent); // Obliczenie rozmiaru tablicy agentow w bajtach

float *host_results = NULL, *device_results = NULL; // Wskaznik do tablic z wynikami
size_t size_results = N * sizeof(float); // Obliczenie rozmiarow tablicy wynikow w bajtach

// Przydzial pamieci
host_agents = (Agent *)malloc(size_agents); // Przydzial pamieci dla agentow na hoscie
cudaMalloc((void **) &device_agents, size_agents); // Przydzial pamieci dla agentow na karcie
host_results = (float *)malloc(size_results); //Przydzial pamieci na wyniki na hoscie
cudaMalloc((void **) &device_results, size_results); //Przydzial pamieci na wyniki na karcie

// Wpisanie wartosci poczatkowych do tablicy
for (int i=0; i < N; i++)
{
host_agents[i].x = (float)(rand()%100);
host_agents[i].y = (float)(rand()%100);
}

//Skopiowanie tablicy z hosta do karty
cudaMemcpy(device_agents, host_agents, size_agents, cudaMemcpyHostToDevice);
// Wywolanie kernela
int block_size = 4;
int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);

Fitness <<< n_blocks, block_size >>> (device_agents, N, device_results);

// Odebranie wyników z karty i zapisanie ich w tablicy na hoscie
cudaMemcpy(host_results, device_results, size_results, cudaMemcpyDeviceToHost);

// Wypisz wyniki
for (int i = 0; i < N; i++)
printf("%05d: f(%04.4f, %04.4f) = %04.4f;\n", i, host_agents[i].x, host_agents[i].y, host_results[i]);
// Cleanup
free(host_agents);
free(host_results);
cudaFree(device_agents);
cudaFree(device_results);
system("pause");
}

wtorek, 20 października 2009

CUDA NVidia

Niedawno otrzymałem zadanie wykorzystania CUDA w algorytmie optymalizacji polegającym na przeszukiwaniu przestrzeni rozwiązań za pomocą "roju". Zaletą CUDA jest wykorzystanie procesora graficznego do przeprowadzania równoległych obliczeń na wielu zmiennych (przynajmniej w tej chwili tak to rozumiem).
Pomyślałem, że mogę wykonać ten projekt na Ubuntu 8.04.
Udałem się na stronę NVIDIA, gdzie w tej chwili znajdują się sterowniki w wersji 185 oraz CUDA Toolkit 2.2.
Niestety poniosłem porażkę - nie umiałem zainstalować sterowników. Po każdej instalacji Xy wstawały w trybie niskiej jakości grafiki. Doczytałem w internecie, że envyNG instaluje sterowniki NVIDIA wraz z biblioteką libcuda. Za pomocą tego programu udało mi się zainstalować sterowniki w wersji 173. Następnie zainstalowałem CUDA Toolkit oraz CUDA SDK. Zacząłem budować przykładowe projekty. Po wykonaniu make dostałem informację, że nie można znaleźć bibliotek libglut, więc zainstalowałem ją z repozytorium - `sudo apt-get install libglut3 libglut3-dev`. Następnym problemem okazał się projekt threadMigration - tutaj po prostu zmieniłem nazwę Makefile na Makefile-renamed i tego projektu nie uda mi się odpalić. W końcu wszystkie projekty zbudowane i wchodzę do katalogu /home/daniel/NVIDIA_CUDA_SDK/bin/linux/release/. Tutaj najpierw uruchamiam ./deviceQuery z następującym rezultatem:
CUDA Device Query (Runtime API) version (CUDART static linking)
There is 1 device supporting CUDA

Device 0: "GeForce 8400M GS"
CUDA Capability Major revision number: 1
CUDA Capability Minor revision number: 1
Total amount of global memory: 267714560 bytes
Number of multiprocessors: 16
Number of cores: 128
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 16384 bytes
Total number of registers available per block: 8192
Warp size: 32
Maximum number of threads per block: 512
Maximum sizes of each dimension of a block: 512 x 512 x 64
Maximum sizes of each dimension of a grid: 65535 x 65535 x 1
Maximum memory pitch: 262144 bytes
Texture alignment: 256 bytes
Clock rate: 0.80 GHz
Concurrent copy and execution: Yes
Run time limit on kernels: No
Integrated: Yes
Support host page-locked memory mapping: Yes
Compute mode: Default (multiple host threads can use this device simultaneously)

Test PASSED

Press ENTER to exit...

Co oznacza, że mam urządzenie "CUDA Capable". Następny test na porozumiewanie się urządzenia z systemem: ./bandwidthTest
Running on......
device 0:GeForce 8400M GS
Quick Mode
Host to Device Bandwidth for Pageable memory
.
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 856.9

Quick Mode
Device to Host Bandwidth for Pageable memory
.
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 895.4

Quick Mode
Device to Device Bandwidth
.
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 4214.5

&&&& Test PASSED

Press ENTER to exit...


Zatem teoretycznie wszystko dobrze, ale nie mogę uruchomić żadnego przykładu. Najczęstszym błędem jest "cudaSafeCall() Runtime API error in file , line 51 : feature is not yet implemented." Wynika z tego, że CUDA 2.2 nie da rady działać na obecnych sterownikach i będę musiał walczyć z tymi ściągniętymi ze strony NVIDIA.

EDIT:
Z ostatniej chwili: udało mi się zainstalować sterowniki na XP. Pomogło wybranie sterowników dla notebooków. Pozwala to wykorzystywać CUDA w wersji 2.2.