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:
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:
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:
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.Podstawowa implementacja OpenCL nie odwzorowuje kolejek logicznych na kolejki fizyczne i decyduje tylko, gdzie umieścić obiekty w środowisku wykonawczym.
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!