Jak skutecznie gromadzić dane z wątków w CUDA?

Question

Mam aplikację, która rozwiązuje system równań w CUDA, wiem na pewno, że każdy wątek może znaleźć do 4 rozwiązań, ale jak mogę skopiować następnie z powrotem do hosta?

Mijam ogromną tablicę z wystarczającą ilością miejsca dla wszystkich wątków, przechowującą 4 rozwiązania (4 podwójne dla każdego rozwiązania), i drugą z liczbą rozwiązań na wątek, ale to naiwne rozwiązanie i jest obecnym wąskim gardłem mojego jądra.

Bardzo lubię to optymalizować. Głównym problemem jest łączna liczba różnych rozwiązań na wątek w jednej macierzy.

Answer 1

Funkcja, której szukasz, nazywa się zagęszczaniem strumienia.

Prawdopodobnie musisz dostarczyć tablicę, która zawiera miejsce dla 4 rozwiązań na wątek, ponieważ próba bezpośredniego zapisania wyników w postaci zwartej może stworzyć tak wiele zależności między wątkami, że wydajność uzyskana dzięki możliwości kopiowania mniej danych z powrotem do hosta jest traconych przez dłuższy czas wykonywania jądra. Wyjątkiem jest to, że prawie wszystkie wątki nie znajdują rozwiązań. W takim przypadku możesz użyć operacji atomowej do utrzymania indeksu w tablicy. Tak więc, dla każdego znalezionego rozwiązania, przechowujesz je w tablicy w indeksie, a następnie użyjesz operacji atomowej do zwiększenia indeksu. Myślę, że byłoby bezpiecznie użyć atomicAdd() do tego. Przed zapisaniem wyniku wątek użyłby metody atomicAdd() w celu zwiększenia indeksu o jeden. atomicAdd() zwraca starą wartość, a wątek może przechowywać wynik, używając starej wartości jako indeksu.

Jednak biorąc pod uwagę częstszą sytuację, w której występuje spora liczba wyników, najlepszym rozwiązaniem będzie wykonanie operacji kompaktowania jako oddzielnego kroku. Jednym ze sposobów na to jest thrust::copy_if. Zobacz this question, aby uzyskać więcej tła.