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");
}