bloki, wątki, warpSize

Question

Było wiele dyskusji o tym, jak wybrać blok blokowy #blocks &, ale nadal czegoś brakuje. Wielu moich obaw rozwiązania tego pytanie: How CUDA Blocks/Warps/Threads map onto CUDA Cores? (Dla uproszczenia dyskusji, nie jest wystarczająco perThread & pamięć perBlock limity pamięci nie są problemem tutaj..)

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal); 

1) Aby utrzymać SM tak zajęty, jak to możliwe, Powinienem ustawić nThreads na wielokrotność warpSize. Prawdziwe?

2) SM może wykonywać tylko jedno jądro naraz. To wszystko, że HWC tego SM wykonuje tylko jądro A. (Nie niektóre HWcores z jądrem A, podczas gdy inne działają z jądrem B.) Więc jeśli mam tylko jeden wątek do uruchomienia, "marnuję" inne HWcores. Prawdziwe?

3) Jeśli problemy z rozkładem wątku działają w jednostkach warpSize (32 wątki), a każdy SM ma 32 HWcores, to SM będzie w pełni wykorzystany. Co się dzieje, gdy SM ma 48 HWcores? Jak mogę w pełni wykorzystać 48 rdzeni, gdy program planujący wydaje pracę w kawałkach 32? (Jeśli poprzedni akapit jest prawdziwy, czy nie byłoby lepiej, gdyby program planujący wydał pracę w jednostkach o rozmiarze HWcore?)

4) Wygląda na to, że program do szeregowania zadań kolejkuje do 2 zadań naraz. Tak więc, gdy aktualnie uruchomione jądro zawiesza się lub blokuje, drugie jądro jest zamieniane. (Nie jest jasne, ale domyślam się, że kolejka tutaj jest większa niż 2 jądra.) Czy to prawda?

5) Jeśli moja HW ma górny limit 512 wątków na blok (nThreadsMax), to nie znaczy, że jądro z 512 wątkami będzie działało najszybciej na jednym bloku. (Znowu mem nie jest problemem.) Jest duża szansa, że ​​osiągnę lepszą wydajność, jeśli rozproszę jądro o długości 512 na wiele bloków, a nie tylko jeden. Blok jest wykonywany na jednym lub wielu SM. Prawdziwe?

5a) Myślę, że im mniej, tym lepiej, ale czy ma znaczenie, jak małe zrobię nBlocks? Pytanie brzmi, jak wybrać wartość nBlocks, która jest przyzwoita? (Niekoniecznie optymalny). Czy istnieje matematyczne podejście do wybierania numeru nBlocks, czy jest to po prostu test-n-err.

Answer 1

Pozwól mi spróbować odpowiedzieć na twoje pytania jeden po drugim.

1. To prawda.
2. Co dokładnie oznacza "HWcores"? Pierwsza część twojego oświadczenia jest poprawna.

3. Zgodnie z NVIDIA Fermi Compute Architecture Whitepaper: "SM planuje wątki w grupach 32 równoległych wątków zwanych osnowami.Każdy SM zawiera dwa układy dopasowania warp i dwie jednostki wysyłania instrukcji, pozwalając na jednoczesne wysyłanie i wykonywanie dwóch osnów. dwie osnowy i wydaje jedną instrukcję z każdej osnowy na grupę szesnastu rdzeni, szesnaście jednostek ładunkowych/magazynowych lub cztery jednostki SFU. Ponieważ osnowy są wykonywane niezależnie, program planujący Fermiego nie musi sprawdzać zależności w strumieniu instrukcji ".

   Co więcej, stwierdza: "Program planujący kwadratu Keplera wybiera cztery osnowy i dwie niezależne instrukcje na osnowę mogą być wysyłane w każdym cyklu."

   Rdzenie "nadmiarowe" są używane do planowania więcej niż jednej osnowy naraz.

4. Harmonogram osnowy planuje wypaczenia tego samego jądra, a nie różne jąder.

5. Nie do końca prawda: każdy blok jest zablokowany w jednym SM, ponieważ tam znajduje się jego pamięć współdzielona.
6. To trudny problem i zależy od sposobu wdrożenia jądra. Możesz rzucić okiem na seminarium nVidia Better Performance at Lower Occupancy autorstwa Wasilija Wołkowa, które wyjaśnia niektóre z ważniejszych kwestii. Przede wszystkim jednak proponuję wybrać liczbę wątków, aby poprawić obłożenie, korzystając z CUDA Occupancy Calculator.

Answer 2

1) Tak.

2) Urządzenia CC 2.0 - 3.0 mogą jednocześnie pracować z maksymalnie 16 siatkami. Każdy SM jest ograniczony do 8 bloków, więc aby osiągnąć pełną współbieżność, urządzenie musi mieć co najmniej 2 SM.

3) Tak, planiści osnowy wybierają i wydają odkształcenia w czasie. Zapomnij o koncepcji rdzeni CUDA, które są nieistotne. Aby ukryć opóźnienie, musisz mieć wysoki poziom równoległości na poziomie instrukcji lub wysokie obłożenie. Zalecane jest> 25% dla CC 1.x i> 50% dla CC> = 2,0. Ogólnie CC 3.0 wymaga większego obłożenia niż urządzenia 2.0 z powodu podwojenia harmonogramów, ale tylko o 33% wzrostu osnów na SM. Eksperyment Nsight VSE Issue Efficiency to najlepszy sposób na ustalenie, czy masz wystarczającą liczbę osnów, aby ukryć instrukcje i opóźnienia w pamięci. Niestety, Visual Profiler nie ma tego pomiaru.

4) Algorytm harmonogramu osnowy nie jest dokumentowany; nie bierze jednak pod uwagę siatki, z której powstał blok wątków. W przypadku urządzeń CC 2.x i 3.0 dystrybutor roboczy CUDA rozprowadzi wszystkie bloki z sieci przed dystrybucją bloków z następnej sieci; jednak nie gwarantuje tego model programowania.

5) Aby utrzymywać SM zajętym, musisz mieć wystarczającą liczbę bloków do napełnienia urządzenia. Po tym chcesz mieć pewność, że masz wystarczającą ilość osnów, aby osiągnąć rozsądne obłożenie. Są zarówno plusy, jak i minusy używania dużych bloków wątków. Bloki dużych wątków generalnie korzystają z mniejszej pamięci podręcznej instrukcji i mają mniejszy ślad w pamięci podręcznej; jednak duże bloki wątków utkną w syncthreadach (SM może stać się mniej wydajnym, ponieważ jest mniej osnów do wyboru) i mają tendencję do utrzymywania instrukcji wykonujących na podobnych jednostkach wykonawczych. Zalecam wypróbowanie 128 lub 256 wątków na blok wątku, aby rozpocząć. Istnieją dobre powody dla większych i mniejszych bloków nici. 5a) Użyj kalkulatora obłożenia. Pobranie zbyt dużego rozmiaru bloku wątków często powoduje, że rejestry są ograniczone. Pobranie zbyt małego fragmentu bloku wątków może ograniczyć dostęp do pamięci współużytkowanej lub do 8 bloków na limit SM.