Optimizacija CUDA cjevovoda za obradu slika: od bugova do benchmarkova

  • PublicPublic
  • 02 Sep, 2026
Optimizacija CUDA cjevovoda za obradu slika: od bugova do benchmarkova

U ovom praktičnom primjeru CUDA obrađuje tok crvenih, zelenih i plavih slika: podaci se najprije kopiraju s CPU-a na GPU, pretvaraju iz RGB-a u nijanse sive, zatim se svaka pločica veličine 32 puta 32 piksela sortira kako bi se mogla izračunati njezina medijana, a na kraju se medijane kopiraju natrag na domaćina. U početnoj implementaciji definirana su dva kernela — computeRGBToGray za pretvorbu boje i computeMedian<TILE_WIDTH, HISTO_SIZE> za medijanu po svakoj pločici, — a main dodjeljuje memoriju, pokreće kernele i oslobađa resurse unutar OpenMP paralelne petlje nad tri slike.

U kernelu medijane krije se suptilan, ali ozbiljan nedostatak. Upis u dijeljenu memoriju tile[index] = d_image_gray[index] pogrešno koristi globalni indeks za upis u dijeljenu memoriju, koja je ograničena na jedan blok niti. Pokretanje binarne datoteke daje pogrešku «illegal memory access», a compute-sanitizer točno ukazuje na uzrok: upis izvan granica __shared__ za jedan bajt u retku 55, zabilježen za nit (0,3,0) u bloku (20,0,0).

Slika 1. Prvi dio cjevovoda za obradu slika. Najprije kopirajte RGB slike s domaćina na uređaj, zatim ih pretvorite iz RGB-a u nijanse sive
Slika 1. Prvi dio cjevovoda za obradu slika. Najprije kopirajte RGB slike s domaćina na uređaj, zatim ih pretvorite iz RGB-a u nijanse sive

Kako bi se spriječile slične pogreške s indeksiranjem, NVIDIA je u CCCL-u predstavila novi API za pokretanje. Kerneli se pokreću putem cuda::make_config, cuda::block_dims, cuda::grid_dims i cuda::launch, zatim kernel prima konfiguraciju kao prvi parametar i koristi cuda::gpu_thread.index(cuda::grid, config) i cuda::gpu_thread.index(cuda::block, config), kako bi odvojeno dohvaćao globalni i blokovski indeks. Sirovi pokazivači također se mogu zamijeniti s cuda::std::span, cuda::std::mdspan i cuda::shared_memory_mdspan, koji u načinu otklanjanja pogrešaka pokreću assertove pri pristupu izvan granica niza.

Slika 2. Drugi dio cjevovoda za obradu slika: podijelite svaku sliku u nijansama sive na pločice od 32 puta 32 piksela, sortirajte piksele u svakoj pločici kako biste odabrali njezinu medijanu, zatim kopirajte vrijednosti medijane iz d
Slika 2. Drugi dio cjevovoda za obradu slika: podijelite svaku sliku u nijansama sive na pločice od 32 puta 32 piksela, sortirajte piksele u svakoj pločici kako biste odabrali njezinu medijanu, zatim kopirajte vrijednosti medijane iz d

Kada kod postane bez pogrešaka, mjeri se pomoću Nsight Systems. Označavanje koda NVTX rasponima — poput nvtx3::scoped_range i nvtxRangePushA/nvtxRangePop — čini vremensku crtu preglednijom i pokazuje da kernel za izračun medijane dominira u profilu: zauzima oko 2,1 sekunde po slici i odgovoran je za 98,5% vremena GPU-a, dok operacije s memorijom zauzimaju samo 1,5%.

Slika 3. Početna Nsight Systems vremenska crta s NVTX oznakama. Kerneli za izračun medijane dominiraju u profilu, čineći gotovo svu aktivnost GPU-a i najveći dio 6,8-sekundne obrade slika
Slika 3. Početna Nsight Systems vremenska crta s NVTX oznakama. Kerneli za izračun medijane dominiraju u profilu, čineći gotovo svu aktivnost GPU-a i najveći dio 6,8-sekundne obrade slika

Sa strane CPU-a cijela faza izračuna nad slikom traje 6,8 sekundi, a gotovo sve to vrijeme odlazi na kernele medijane. Takva vremenska crta daje programeru jasan cilj za sljedeće korake optimizacije: izračun medijane najskuplji je dio cjevovoda, a kasnije faze prikaza usmjerene su upravo na njegovo poboljšanje.