Klasyczny wzór indeksu

Cały sens kernela CUDA polega na tym, że każdy wątek obsługuje jeden fragment danych. Aby to osiągnąć, każdy wątek potrzebuje unikalnego indeksu globalnego.

Wewnątrz bloku threadIdx.x przyjmuje tylko wartości od 0 do blockDim minus jeden. Wiele bloków używa ponownie tych samych małych liczb, więc nie mogą one być indeksem końcowym.

Wyobraź sobie siatkę jako bloki ułożone jeden za drugim. Każdy blok odpowiada za spójny fragment tablicy, a blockIdx.x wskazuje, w którym fragmencie się znajdujesz.

blockDim.x to liczba wątków w bloku. Określa szerokość każdego fragmentu, więc skaluje przesunięcie bloku do właściwej pozycji.

Połącz te trzy elementy: pomiń wcześniejsze bloki, a następnie dodaj swoją pozycję w bieżącym bloku. Ta jedna linia nadaje każdemu wątkowi unikalny indeks.

Przy 4 wątkach na blok blok 0 obejmuje zakres od 0 do 3, blok 1 — od 4 do 7, a blok 2 — od 8 do 11. Przesunięcia nigdy się nie nakładają.

Wątek 2 w bloku 3 przy blockDim równym 256 trafia na pozycję 3 razy 256 plus 2, czyli 770. To jego globalna pozycja w danych.

Gdy masz już i, traktujesz je jako indeks tablicy. Każdy wątek odczytuje i zapisuje wyłącznie własny element, bez nakładania się operacji.

Uruchom n wątków, a wzór wygeneruje każdą wartość od 0 do n minus 1 dokładnie raz. To idealne odwzorowanie jeden do jednego całej tablicy.

Zawsze najpierw wykonuj mnożenie, a dopiero potem dodawanie. blockIdx.x * blockDim.x to początek Twojego fragmentu, a threadIdx.x określa pozycję w jego obrębie.

Ta sama idea rozszerza się na 2D i 3D z użyciem składowych .y i .z, ale w przypadku płaskich tablic wystarczy Ci wzór z .x. 🚀

Poznałeś wzór używany przez każdy kernel: blockIdx.x * blockDim.x + threadIdx.x. Przydziela on każdemu wątkowi jedno unikalne miejsce w tablicy. 🎉