Iterowanie po fazach kafelków

Pełny wiersz pomnożony przez pełną kolumnę jest zbyt duży dla jednego kafelka. Dlatego długą sumę dzieli się na fragmenty o szerokości TILE, po jednym fragmencie na każdą fazę.

Jeśli macierze mają szerokość N, a kafelki szerokość TILE, potrzeba N / TILE faz, aby objąć cały wymiar wewnętrzny.

Umieść kroki wczytywania, synchronizacji i obliczeń w pętli po phase. Każde wykonanie przesuwa okno kafelka dalej wzdłuż wiersza A i kolumny B.

Lokalna zmienna sum znajduje się poza pętlą, więc jej wartość stale rośnie. Każda faza dodaje swój fragment iloczynu skalarnego do bieżącej sumy.

W każdej fazie kolumna odczytywana z A i wiersz odczytywany z B przesuwają się o phase * TILE. W ten sposób przesuwa się okno.

Tak jak wcześniej, wywołaj __syncthreads() po wczytaniu danych, aby kafelek był kompletny, zanim którykolwiek wątek zacznie na nim obliczenia.

Pętla wewnętrzna dodaje TILE iloczynów do sumy, korzystając wyłącznie ze świeżo wczytanego kafelka. Wnosi ona jeden fragment do końcowego iloczynu skalarnego.

Każdą fazę zakończ kolejnym wywołaniem __syncthreads(), aby żaden wątek nie nadpisał kafelka, gdy wolniejszy wątek nadal go odczytuje. 🚧

Jedna bariera chroni odczyty wykonywane po wczytaniu, a druga chroni wczytywanie wykonywane po odczytach. Razem utrzymują każdy wątek w synchronizacji podczas kolejnych faz.

Po zakończeniu wszystkich faz bieżąca suma jest pełnym iloczynem skalarnym. Zapisz ją do C raz, zabezpieczając zapis sprawdzeniem zakresu.

Gdy N nie jest całkowitą wielokrotnością TILE, dla elementów spoza zakresu wczytaj zero, aby dodatkowe iloczyny nic nie dodały do sumy.

Zastanów się, gdzie znajduje się bieżąca suma podczas pętli faz.

Wykonali Państwo pętlę po fazach, przesuwając kafelki, synchronizując wątki dwukrotnie i dodając sumy częściowe do jednej wartości. Teraz zmierzą Państwo, o ile rozwiązanie działa szybciej. 📈