CUDA - obliczenia na układach graficznych

Materiały "najnowsze"

Materiały historyczne, ale możliwe, że pożyteczne

Co się działo na zajęciach

  1. Zajęcia 1: 2014-02-17
    • Ogólne wprowadzenie do obliczeń na kartach.
    • Praca z materiałami źródłowymi
      • NVIDIA Fermi Architecture
      • NVIDIA Kepler Architecture
      • CUDA, Supercomputing for the Masses: Part 1 - 2 (introduction to CUDA and first CUDA program)
      • Analiza prgramu deviceQuery
        C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5\bin\win32\Release
        
        Powyższy link prowadzi do wersji wykonywalnej, źródło jest w lokalizacji
        C:\ProgramData\NVIDIA Corporation\?????
        
  2. Zajęcia 2: 2014-02-17
    • Podsumowanie pracy z materiałami źródłowymi z poprzednich zajęć.
    • Piszemy pierwszy program na CUDA: program dodaje 1 (lub 1.0 - warto sprawdzić różnice w wydajności) do elementów (małego) wektora.
      Przydatne rzeczy:
      • CUDA C Programming Guide
        • 3.2.9. Error Checking
        • CUDA, Supercomputing for the Masses: Part 3 Error handling and global memory performance limitations
        • // Simple utility function to check for CUDA runtime errors
          void checkCUDAError(const char* msg);
          
          
          void checkCUDAError(const char *msg)
          {
              cudaError_t err = cudaGetLastError();
              if( cudaSuccess != err) 
              {
                  fprintf(stderr, "Cuda error: %s: %s.\n", msg, 
                                            cudaGetErrorString( err) );
                  exit(EXIT_FAILURE);
              }                         
          }
          
      • Pomiar czasu na CUDA-ch
        	
        // create the timer	
        cudaEvent_t start, stop;
        float elapsedTime;
        
        // start a timer
        cudaEventCreate( &start );
        cudaEventCreate( &stop );
        cudaEventRecord( start, 0 );
        
        // do something
        
        // stop a timer and collect performance data
        cudaEventRecord( stop, 0 );
        cudaEventSynchronize( stop );
        cudaEventElapsedTime( &elapsedTime, start, stop );
        printf( "Time taken: %3.1f ms\n", elapsedTime );
        
    Co po tych zajęciach powinno się znać / wiedzieć / umieć
    • Krótka historia powstania GPU i obliczeń bazujacych na GPU.
    • Czym różni się CPU od GPU? Czy możliwe jest zastąpienie jednego przez drugie?
    • Najważniejsze koncepcje związane z CUDA (kernel, wątek, blok, grid, hierarchia pamięci).
    • Architektura Fermi ze szczególnym uwzględnieniem architektury SM.
    • Architektura Kepler - istotne różnice w porównaniu do Fermi.
    • Napisać najprostszy program (wykorzystując schemat: move data to CUDA-enabled device(s), perform a calculation and retrieve result) i wyjaśnić rolę jego poszczególnych elementów (podział wątków na bloki i gridy, sposób wywołania kernela).
    Praca z materiałami źródłowymi
  3. Zajęcia 3: 2014-03-03
    • Podsumowanie pracy z materiałami źródłowymi z poprzednich zajęć.
    • Programy do zaimplementowania
      1. Proste dodawanie +1 dla malego wektora (mieszczącego się w jednym bloku) - zadanie z poprzednich zajęć, jeśli ktoś tego jeszcze nie zrobił.
      2. Dodawanie +1 dla dużego wektora (wymagającego użycia wielu bloków), ale mieszczącego się w pamięci - kontynuacja zadania z poprzednuch zajęć.
      3. Dodawanie +1 dla wektora który w całości nie mieści się w pamięci karty
      4. Odwracanie wektora na pamięci globalnej.
      5. Odwracanie wektora na pamięci współdzielonej.
      6. Pamięć współdzielona - wywoływanie konfliktów.
      7. Uruchomienie kerneli, które nie używają wszystkich rdzeni.
      UWAGA! Każdy z uczestników zajęć wysyła do mnie implementacje powyższych zadań. Proszę je przesyłać (najlepiej same źródła) w postaci spakowanej (zip, bardzo proszę aby nie był to rar) stosując dla przesyłanego pliku następującą konwencję nazewniczą
      1. Nie używać dużych liter i polskich znaków.
      2. Nazwa pliku powinna pasować do schematu: nazwisko_imie_nrZajec_nrProjektu.zip
      3. nrZajec i nrProjektu muszą być dwucyfrowe, więc w razie konieczności nalży pojedyńczą cyfrę poprzedzić zerem.
  4. Zajęcia 4: 2014-03-10
    • Kończymy programy z poprzednich zajęć
    • Podsumowanie pracy z zajęć 3.
    • Praca z materiałami źródłowymi
    • Programy do zaimplementowania
      1. Napisać program(y), który pokaże działanie funkcji __syncthreads() oraz dowolnej innej z rodziny __syncthreads.
      2. Obliczenie iloczynu skalarnego metodą redukcji. Uwaga! Metoda redukcji nie ma nic wspólnego z matematyką :) tylko z metodą postępownia którą ilustruje poniższy obrazek

        Warto też przeczytać Optimizing Parallel Reduction in CUDA.
      3. Napisać program, który pokaże działanie dwóch funkcji z rodziny __threadfence.
    Co po tych zajęciach (zajęcia 3-4) powinno się znać / wiedzieć / umieć
    • Różnice w wydajności pomiędzy pamięcią globalna a współdzieloną.
    • Wpływ na wydajność operacji na pamięci w zależności od wzorca dostępu do niej.
    • Metody synchronizacji wątków.
    • Metody ,,synchronizacji'' operacji na pamięci.
  5. Zajęcia 5: 2014-03-17
    • Kończymy programy z poprzednich zajęć
    • Podsumowanie pracy z zajęć 4.
    • Praca z materiałami źródłowymi
    • Programy do zaimplementowania
      1. Napisać poprawnie działający program obliczający histogram z użyciem funkcji atomowych i bez nich.
      2. Napisać program(y), który pokaże działanie funkcji __all(), __any() oraz __ballot() z rodziny Warp Vote Functions.
    Co po tych zajęciach powinno się znać / wiedzieć / umieć
    • Umieć podać przykład problemu wraz z rozwiązaniem (inny niż omawiany na zajęciach), który wymaga użycia funcji atomowych.
    • Umieć opisać działanie funkcji z rodziny Warp Vote Functions.
  6. Zajęcia 6: 2014-03-24
    • Nic nowego - kończymy programy z poprzednich zajęć
    • Ciekawa dyskusja z p. Pawłem Jarzyńskim na temat opcji kompilacji -arch i -code (podsumowanie)
  7. Zajęcia 7: 2014-03-31
  8. Zajęcia 8: 2014-04-07
    • Nic nowego - kończymy programy z poprzednich zajęć
  9. Zajęcia 9: 2014-04-14
    • Kończymy programy z poprzednich zajęć
    • Programy do zaimplementowania
      1. Program wyszukujący zadanego podciągu w ciągu. Program należy porównać z dowolnym (wydajnym) algorytmem działającym na CPU.
      2. Program wyszukujący najmniejszego i największego elementu w ciągu liczb. W przypadku tego algorytmu można przyjąć pewne ograniczenia na liczby w ciągu jeśli pozwoli to znacząco podnieść wydajność Państwa algorytmu.
    • Ciekawostka: Programming Guide: B.14. Warp Shuffle Functions i mój test
  10. Zajęcia 10: 2014-04-28
  11. Zajęcia 11: 2014-05-05
  12. Zajęcia 12: 2014-05-12
  13. Zajęcia 13: 2014-05-19
    • Kończymy programy z poprzednich zajęć
    • Programy do zaimplementowania
      1. Napisać program sprawdzający czy w grafie istnieje cykl Hamiltona.
  14. Zajęcia 14: 2014-05-26