• W liniach 35–40 uzyskane zostały obiekty platformy platform oraz urzą- dzenia device...

Pokaż mi serce nie opętane zwodniczymi marzeniami, a pokażę ci człowieka szczęśliwego.

Na aktualne zostało wybrane pierwsze urządzenie typu
CL_DEVICE_TYPE_GPU zainstalowane w systemie. Z tym urządzeniem został
skojarzony kontekst context, a następnie została dla niego stworzona
kolejka poleceń cmdQueue.
•
W liniach 42–47 wczytywany jest z pliku matadd.cl kod źródłowy kernela
(przedstawiony na listingu 2.7). Na podstawie tego kodu tworzony jest
w linii 47 obiekt programu hProgram, który jest następnie budowany w
linii 49.
•
W liniach 51–58 za pomocą funkcji clCreateBuffer() utworzone zostały 3
obiekty pamięciowe matA, matB, matC, które będą reprezentować wartości
macierzy po stronie urządzenia OpenCL.
•
W liniach 60–63 dane z macierzy hosta są kopiowane do pamięci GPU.
•
W liniach 65–69 zostały wyspecyfikowane argumenty wywołania funkcji
kernela.
•
W liniach 71–74 następuje właściwe zakolejkowanie funkcji kernela za po-
mocą funkcji clEnqueueNDRangeKernel(). W tym przypadku globalna ilość
56
2. Architektura środowisk CUDA i OpenCL
work-items GLOBAL_WS została ustawiona na wielkość macierzy docelowej
matC a lokalna LOCAL_WS na 16 × 16.
Na koniec pozostaje jeszcze przedstawienie samej funkcji rdzenia liczącej
sumy poszczególnych elementów macierzy:
Listing 2.7. OpenCL – Dodawanie macierzy – funkcja rdzenia.
1
__kernel void matadd ( __global float * matDest ,
2
__global float * matA ,
3
__global float * matB ,
4
int len )
5
{
6
int i = get_global_id (0) + get_global_id (1) *
7
get_global_size (0) ;
8
if (i < len )
9
matDest [i] = matA [i] + matB [i];
10
}
2.5. Pomiar czasu za pomocą zdarzeń GPU
Dotychczas stosowanym sposobem pomiaru czasu była dosyć nieprecy-
zyjna metoda oparta na pobieraniu czasu systemowego, którego dokładność
wynosiła ok. 1 milisekundy. Ta dokładność w dużej mierze zależy od pro-
gramowanego sprzętu i systemu operacyjnego. Co więcej, problematyczny
może okazać się pomiar czasu funkcji asynchronicznych wykonywanych na
GPU. Rozwiązaniem tego problemu może być użycie dedykowanych funk-
cji pomiaru czasu zdefiniowanych w obu środowiskach CUDA/OpenCL. W
obu przypadkach pomiar czasu jest oparty na tzw. obiektach zdarzeniowych
(ang. event obects). Czas jest w tym przypadku odmierzany przez urządzenie
obliczeniowe (GPU) z dokładnością do nanosekund.
Obiekt zdarzeniowy w CUDA cudaEvent_t event jest tworzony za pomocą
AD funkcji konstruktora:
UC
cudaError_t cudaEventCreate ( cudaEvent_t * event )
i niszczony za pomocą funkcji:
cudaError_t cudaEventDestroy ( cudaEvent_t event )
Takie obiekty zdarzeniowe mogą zostać wykorzystane do pomiaru czasu
przez ich rejestrację w danym strumieniu za pomocą funkcji:
cudaError_t cudaEventRecord ( cudaEvent_t event ,
2.5. Pomiar czasu za pomocą zdarzeń GPU
57
cudaStream_t stream = 0)
Tak zarejestrowane zdarzenie zostanie uznane za zakończone, gdy zostaną
wykonane wszystkie polecenia z danego strumienia poprzedzające wywoła-
nie funkcji rejestrującej zdarzenie. Zakończone zdarzenia mogą posłużyć do
pomiaru czasu dzięki funkcji:
cudaError_t cudaEventElapsedTime ( float * ms , cudaEvent_t start ,
cudaEvent_t end )
określającej, z bardzo dużą dokładnością, czas jaki upłyną po stronie GPU
pomiędzy tymi dwoma zdarzeniami. Dokładnie, funkcja cudaElapsedTime()
zwraca czas jaki upłyną pomiędzy zakończeniem zdarzenia end a zakończe-
niem zdarzenia start.
Poniżej znajduje się fragment prostego kodu obrazującego sposób po-
miaru czasu metodą zdarzeniową:
Listing 2.8. CUDA – Metoda pomiaru czasu za pomocą zdarzeń.
32
float time ;
33
cudaEvent_t start , end ;
34
35
cudaEventCreate (& start )
36
cudaEventCreate (& end )
37
38
cudaEventRecord ( start , 0) ;
39
40
cudaMemcpy (...)
41
kernel < < <... > > >(...) ;
42
cudaMemcpy (...)
43
44
cudaEventRecord (end , 0) ;
45
cudaEventSynchronize ( end );
46
cudaEventElapsedTime (& time , start , end );
Przed
samym
pomiarem
czasu
została
wywołana
funkcja
cudaEventSynchronize(cudaEvent_t). Jest to jedna z wielu metod syn-
chronizacji strumienia CUDA po stronie hosta, która blokuje aktualny
wątek hosta, aż do czasu wykonania wszystkich poleceń w danym strumieniu
poprzedzających podane w parametrze zdarzenie.
W OpenCL, aby możliwy był pomiar czasu GPU, należy włączyć opcję
Ope
profilowania kolejki rozkazów podczas jej tworzenia:
nCL
cl_command_queue queue = clCreateCommandQueue ( context , devices ,
CL_QUEUE_PROFILING_ENABLE , 0) ;
58
2. Architektura środowisk CUDA i OpenCL
W powyższym przykładzie, w trzecim parametrze została podana flaga
CL_QUEUE_PROFILING_ENABLE włączająca profilowanie. W przypadku istnieją-
cych kolejek rozkazów można włączać lub wyłączać konkretne opcje za po-
mocą funkcji:
cl_int clSetCommandQueueProperty ( cl_command_queue command_queue ,
cl_command_queue_properties properties ,
cl_bool enable ,
cl_command_queue_properties * old_properties )
Listing 2.9 przedstawia sposób wykorzystania obiektów zdarzeniowych
cl_events do pomiaru czasu wykonania kolejki poleceń oraz czasu wykonania
pojedynczego polecenia.
Listing 2.9. OpenCL – Metoda pomiaru czasu za pomocą zdarzeń.
32
cl_event start , end , kernel_ev ;
33
cl_ulong time , time2 ;
34
35
clEnqueueMarker ( queue , & start );
36
37
clEnqueueWriteBuffer (...) ;
38
clSetKernelArg (...) ;
39
clEnqueueNDRangeKernel ( queue , kernel , 1, 0,
40