Na co czeka mój GPU?

11

Piszę program OpenCL do użytku z moim procesorem graficznym AMD Radeon z serii HD 7800. Według przewodnika programowania OpenCL AMD , ta generacja GPU ma dwie kolejki sprzętowe, które mogą działać asynchronicznie.

5.5.6 Kolejka poleceń

W przypadku wysp południowych i późniejszych urządzenia obsługują co najmniej dwie kolejki obliczeniowe sprzętu. Pozwala to aplikacji na zwiększenie przepustowości małych wysyłek dzięki dwóm kolejkom poleceń do asynchronicznego przesyłania i ewentualnie wykonywania. Sprzętowe kolejki obliczeniowe są wybierane w następującej kolejności: pierwsza kolejka = parzyste kolejki poleceń OCL, druga kolejka = nieparzyste kolejki OCL.

Aby to zrobić, utworzyłem dwie osobne kolejki poleceń OpenCL, aby przesłać dane do GPU. Z grubsza program działający w wątku hosta wygląda mniej więcej tak:

static const int kNumQueues = 2;
cl_command_queue default_queue;
cl_command_queue work_queue[kNumQueues];

static const int N = 256;
cl_mem gl_buffers[N];
cl_event finish_events[N];

clEnqueueAcquireGLObjects(default_queue, gl_buffers, N);

int queue_idx = 0;
for (int i = 0; i < N; ++i) {
  cl_command_queue queue = work_queue[queue_idx];

  cl_mem src = clCreateBuffer(CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, ...);

  // Enqueue a few kernels
  cl_mem tmp1 = clCreateBuffer(CL_READ_WRITE);
  clEnqueueNDRangeKernel(kernel1, queue, src, tmp1);

  clEnqueueNDRangeKernel(kernel2, queue, tmp1, tmp1);

  cl_mem tmp2 = clCreateBuffer(CL_READ_WRITE);
  clEnqueueNDRangeKernel(kernel2, queue, tmp1, tmp2);

  clEnqueueNDRangeKernel(kernel3, queue, tmp2, gl_buffer[i], finish_events + i);

  queue_idx = (queue_idx + 1) % kNumQueues;
}

clEnqueueReleaseGLObjects(default_queue, gl_buffers, N);
clWaitForEvents(N, finish_events);

Dzięki kNumQueues = 1tej aplikacji działa ona właściwie tak, jak powinna: gromadzi całą pracę w jednej kolejce poleceń, która następnie jest uruchamiana do końca, a procesor graficzny jest dość zajęty przez cały czas. Widzę to, patrząc na dane wyjściowe profilera CodeXL:

wprowadź opis zdjęcia tutaj

Jednak kiedy ustawiam kNumQueues = 2, spodziewam się tego samego, ale z równomiernym podziałem pracy na dwie kolejki. Jeśli już, oczekuję, że każda kolejka będzie miała te same cechy osobno co jedna kolejka: że zacznie działać sekwencyjnie, aż wszystko zostanie zrobione. Jednak podczas korzystania z dwóch kolejek widzę, że nie cała praca jest podzielona na dwie kolejki sprzętowe:

wprowadź opis zdjęcia tutaj

Na początku pracy GPU kolejkom udaje się uruchamiać niektóre jądra asynchronicznie, chociaż wydaje się, że żadne z nich w pełni nie zajmuje się kolejkami sprzętowymi (chyba że moje zrozumienie jest błędne). Pod koniec pracy GPU wydaje się, że kolejki dodają pracę sekwencyjnie tylko do jednej z kolejek sprzętowych, ale zdarzają się nawet czasy, gdy żadne jądra nie są uruchomione. Co daje? Czy mam jakieś fundamentalne nieporozumienie dotyczące tego, jak powinien się zachowywać środowisko wykonawcze?

Mam kilka teorii, dlaczego tak się dzieje:

  1. Przerywane clCreateBufferwywołania zmuszają procesor GPU do alokacji zasobów urządzenia z puli pamięci współużytkowanej synchronicznie, co wstrzymuje wykonywanie poszczególnych jąder.

  2. Podstawowa implementacja OpenCL nie odwzorowuje kolejek logicznych na kolejki fizyczne i decyduje tylko, gdzie umieścić obiekty w środowisku wykonawczym.

  3. Ponieważ używam obiektów GL, GPU musi synchronizować dostęp do specjalnie przydzielonej pamięci podczas zapisywania.

Czy którekolwiek z tych założeń są prawdziwe? Czy ktoś wie, co może powodować, że GPU czeka w scenariuszu z dwiema kolejkami? Będziemy wdzięczni za wszelkie informacje!

gpgpu opencl Mokosha
źródło
Nie mogę znaleźć, gdzie jest napisane, że w przewodniku programowania są dwie kolejki sprzętowe. Czy możesz opublikować wycenę z dokumentu? Wspomnieć, który rozdział mówi, że są dwie kolejki? Czy w środowisku wykonawczym można zapytać o liczbę kolejek sprzętowych przy użyciu OpenCL?
Andreas
Zaktualizowałem swój post. Mówi o możliwej realizacji, ale jeśli może zrobić kilka, dlaczego nie może zrobić wszystkich? Również środowisko wykonawcze OpenCL nie ma pojęcia o kolejce sprzętowej, więc nie można o to pytać.
Mokosha

Odpowiedzi:

2

Generalnie kolejki obliczeniowe niekoniecznie oznaczają, że możesz teraz wykonywać 2x wysyłki równolegle. Pojedyncza kolejka, która w pełni nasyca jednostki obliczeniowe, będzie miała lepszą przepustowość. Wiele kolejek jest użytecznych, jeśli jedna kolejka zużywa mniej zasobów (pamięć współdzielona lub rejestry), a następnie kolejki dodatkowe mogą się nakładać na tę samą jednostkę obliczeniową.

W przypadku renderowania w czasie rzeczywistym jest to szczególnie ważne w przypadku renderowania w tle, które jest bardzo lekkie dla komputerów / shaderów, ale ciężkie dla sprzętu o stałej funkcji, dzięki czemu program planujący GPU może uruchomić asynchronizację kolejki dodatkowej.

Znalazłem to również w uwagach do wydania. Nie wiem, czy to ten sam problem, ale może być tak, że CodeXL nie jest świetny. Spodziewałbym się, że może nie mieć najlepszego oprzyrządowania, do którego wysyłane są przesyłki.

https://developer.amd.com/wordpress/media/2013/02/AMD_CodeXL_Release_Notes.pdf

W przypadku aplikacji, która wykonuje jednoczesne asynchroniczne przesyłanie danych i wykonywanie jądra, oś czasu pokazana w widoku sesji śledzenia aplikacji nie pokazuje nakładania się tych operacji. Wynika to z faktu, że sterownik i sprzęt wymuszają synchronizację tych operacji podczas profilowania. (333981)

Calvin
źródło