Przeczytałem https ://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-bar zawiera szczegółowe informacje o funkcji synchronizacji PTX.

  1. Mówi, że istnieje 16 „zasobów logicznych bariery”, a za pomocą parametru „a” możesz określić, której bariery użyć. Co to jest zasób logiczny bariery?

  2. Mam fragment kodu z zewnętrznego źródła, o którym wiem, że działa. Jednak nie rozumiem składni używanej w "asm" i co robi "pamięć". Zakładam, że „name” zastępuje „%0”, a „numThreads” zastępuje „%1”, ale co to jest „pamięć” i co robią dwukropki?

    __device__ __forceinline__ void namedBarrierSync(int name, int numThreads) {
    asm volatile("bar.sync %0, %1;" : : "r"(name), "r"(numThreads) : "memory");}
    
  3. W bloku 256 wątków chcę zsynchronizować tylko wątki 64 ~ 127. Czy to możliwe dzięki barrier.sync funkcjonować? ( na przykład, powiedzmy, że mam siatkę 1 bloku, blok 256 wątków. dzielimy blok na 3 gałęzie warunkowe st wątki 0 ~ 63 przechodzą do jądra1, wątki 64 ~ 127 przechodzą do jądra 2, a wątki 128 ~ 255 przejdź do jądra 3. Chcę, aby wątki w jądrze 2 synchronizowały się tylko między sobą.Więc jeśli używam funkcji „namedBarrierSync” zdefiniowanej powyżej: „namedBarrierSync( 1, 64)”.Wtedy synchronizuje tylko wątki 64 ~ 127 lub wątki 0 ~ 63?

  4. Testowałem z poniższym kodem (załóżmy, że gpuAssert jest funkcją sprawdzania błędów zdefiniowaną gdzieś w pliku).

Oto kod:

__global__ void test(int num_threads) 
{
    if (threadIdx.x >= 64 && threadIdx.x < 128) 
    {
        namedBarrierSync(0, num_threads) ;
    }
    __syncthreads();
}

int main(void) 
{
    test<<<1, 1, 256>>>(128);
    gpuAssert(cudaDeviceSynchronize(), __FILE__, __LINE_);
    printf("complete\n");
    return 1;
}
1
Cucumber Mellon U. 7 grudzień 2018, 05:40

1 odpowiedź

Najlepsza odpowiedź
  1. "zasób logiczny bariery" to sprzęt niezbędny do synchronizacji wątków/wypaczeń w bloku wątków (prawdopodobnie liczniki atomowe itp.). Nie musisz znać rzeczywistej implementacji sprzętowej, aby je zaprogramować, wystarczy wiedzieć, że dostępnych jest 16 instancji.
  2. Jak zauważył Robert Crovella w swoim cross-post na forum Nvidii, dokumentacja wbudowanego PTX znajduje się pod adresem https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html.
  3. barrier.sync z nazwaną barierą i liczbą wątków 64 synchronizuje pierwsze dwa wypaczenia docierające do nazwanej bariery (dla możliwości obliczeniowych do 6.x) lub pierwsze 64 wątki docierające do nazwanej bariery (dla możliwości obliczeniowych 7.0 i nowszych ).
  4. Twój test uruchamia tylko jeden wątek (z przydzielonymi do niego 256 bajtami pamięci współdzielonej), co sprawia, że ​​testy instrukcji synchronizacji stają się dyskusyjne. Zamiast tego chcesz uruchomić jądro testowe jako test<<<1, 256>>>(128);.
3
tera 7 grudzień 2018, 13:16