Synchronizowanie za pomocą __syncthreads

Wątki w bloku nie działają krok w krok. Jeden może zakończyć zapis, gdy inny wciąż odczytuje dane, więc nie można zakładać, że dane są już gotowe.

Bariera to miejsce w kodzie, w którym każdy wątek musi czekać, aż dotrą do niego wszystkie pozostałe. Dopiero wtedy blok przechodzi dalej wspólnie.

CUDA udostępnia __syncthreads() jako barierę obejmującą cały blok. Wywołaj ją, a każdy wątek w bloku zatrzyma się do chwili przybycia ostatniego wątku. 🛑

Podstawowy wzorzec jest prosty: każdy wątek ładuje swój fragment do pamięci współdzielonej, następuje synchronizacja, a potem wszyscy bezpiecznie odczytują dane zapisane przez pozostałe wątki.

Bez bariery wątek może odczytać element, zanim sąsiedni wątek go zapisze. To wyścig danych, którego wynikiem są ciche, błędne dane.

Bariera sprawia również, że zapisy do pamięci współdzielonej stają się widoczne dla wszystkich wątków po jej przekroczeniu. Po przejściu przez __syncthreads wszyscy widzą więc świeżo zapisane wartości.

Ścisła zasada brzmi: każdy wątek w bloku musi dotrzeć do tego samego __syncthreads. Jeśli niektóre wątki je pominą, blok może zawiesić się na zawsze.

Nigdy nie umieszczaj bariery wewnątrz instrukcji if, do której wchodzą tylko niektóre wątki. Wątki wybierające drugą ścieżkę nigdy nie dotrą do bariery, a blok zakleszczy się. ⚠️

Ważne ograniczenie: __syncthreads synchronizuje tylko jeden blok. Nie może koordynować wątków z różnych bloków siatki.

W algorytmach kafelkowych często synchronizujesz wątki dwa razy na fazę: raz po załadowaniu kafelka i raz po obliczeniach, przed załadowaniem następnego.

Bariera kosztuje trochę czasu, ponieważ wątki muszą czekać. Używaj jej tam, gdzie wymaga tego poprawność, ale unikaj zbędnych wywołań, które tylko zatrzymują szybsze wątki.

Dowiedziałeś się, że __syncthreads() to bariera obejmująca cały blok, która utrzymuje synchronizację wątków, a umieszczenie jej w rozbieżnych gałęziach może prowadzić do zakleszczenia. Dalej: konflikty banków. 🎯