Savunma tarafında gerçek zamanlı bir tespit ve takip hattı üzerinde çalışıyorum: görünür ve termal video, C++, NVIDIA Jetson üzerinde. Aynı programın içinde hem OpenMP hem CUDA C var. Bana sık sık hangisi daha iyi diye soruluyor ve soru yanlış. İkisi farklı şeyleri paralelleştirir. Asıl beceri, hattın hangi aşamasının hangisine ait olduğuna karar vermek; çünkü ikisi de aynı bütçeden harcar: kare başına elindeki milisaniyeler.
Şekil: sentetik veri üzerinde kendi çizimim.
İki farklı paralellik biçimi
OpenMP, CPU üzerinde paylaşımlı bellekli paralellik. Elinde bir avuç güçlü
çekirdek var. Bir döngüyü #pragma omp parallel for ile işaretliyorsun,
çalışma zamanı iterasyonları çekirdeklere dağıtıyor. Başlamanın maliyeti
neredeyse sıfır. Ve her çekirdek, GPU’nun nefret ettiği şeylerde iyi:
dallanma, işaretçi takibi, her iterasyonun farklı iş yaptığı kod.
CUDA bunun tam tersi biçim. Binlerce hafif iş parçacığı ve hepsi aynı komutu farklı veri üzerinde çalıştırmak istiyor. Piksel başına iş, ders kitabı örneği: yeniden boyutlandırma, renk dönüşümü, filtre çekirdeği. Ama bir warp içindeki iş parçacıkları farklı dallara saparsa serileşirler; dallı kod makineyi boşa harcar.
İkisinin de faturası ancak yük altında görünen ince yazıları var. OpenMP tarafında klasik tuzak false sharing: aynı önbellek satırında yaşayan komşu elemanları güncelleyen iki iş parçacığı, paralel döngünü sessizce bir önbellek ping-pong maçına çevirir. Diğeri aşırı abonelik; OpenMP bütün çekirdekleri kapar, oysa sürecin geri kalanının da birkaç çekirdeğe ihtiyacı vardır. CUDA tarafındaki ince yazı ise başlatma maliyeti ve doluluk: bir çekirdek başlatmanın sabit bir bedeli var, birkaç düzine eleman için çekirdek ateşlemek beş kişilik toplantı için stadyum kiralamak gibi. İki ders de aynı yöne işaret ediyor: aracı işin biçimi seçer, tersi değil.
Taşıma vergisi
CUDA’nın öğrettiği ilk şey: pahalı olan nadiren çekirdeğin kendisidir. Asıl fatura veri taşımaktır. Masaüstünde kare PCIe üzerinden geçer. Jetson’da CPU ile GPU aynı fiziksel belleği paylaşır, bu yardımcı olur; ama özensiz bir hat yine de kareleri sebepsiz yere oradan oraya kopyalar.
Kuralım şu: kareyi GPU’ya bir kez çıkar, ön işleme ve çıkarım boyunca orada tut, geri yalnızca küçük şeyleri getir: kutular, skorlar, takip kimlikleri. Karenin kendisini asla. Bu ayrımı CUDA, CUDA C ve OpenCV CUDA modülü yazısında daha derin anlattım.
Jetson’da bunun bir katmanı daha var. Paylaşılan DRAM, taşımanın bedava olduğu anlamına gelmiyor: varsayılan OpenCV yolu bir Mat’i GpuMat’e yüklerken yine ara kopya kullanır ve her senkronizasyon noktası, CPU ile GPU birbirini beklerken hattı durdurur. Sabitlenmiş ya da yönetilen bellek, bir kez ayrılıp her karede yeniden kullanılan tamponlar ve asenkron akışlar bu acının çoğunu alır. Bende kalan alışkanlık şu: sıcak yoldaki her cudaMemcpy, yaşamayı hak ettiğini kanıtlayana kadar hatadır.
Takip hattında kim nerede kazanır
Takip işinde yerleşim şöyle oturdu. GPU: distorsiyon düzeltme, yeniden boyutlandırma, renk dönüşümü, çıkarım. Milyonlarca piksel üzerinde tekdüze aritmetik; CUDA tam bunun için var. OpenMP’li CPU: nesne başına olan her şey. Eşleştirme, filtre güncellemesi, takip yönetimi, operatör için çizim. Yoğun bir sahnede milyonlarca piksel değil onlarca takip var ve mantık dallanma dolu. GPU’nun orada çiğneyecek bir şeyi yok. Bir CPU çekirdeği o işi yer bitirir.
OpenMP bir de hat tutkalı olarak ekmeğini çıkarıyor: yakalama, işleme ve gösterimi üst üste bindiren paralel bölümler sayesinde aşamalar birbirinin gecikmesini gizliyor. Bu, paralellikten çok bir multithreading hikâyesi; onu ayrıca gerçek zamanlı C++ multithreading yazısında anlattım.
Bir aşamayı yönlendiren üç soru
Yeterince deneme yanılmadan sonra bu karar zevk meselesi olmaktan çıktı, her aşamaya sorduğum üç soruya dönüştü.
Birincisi: iş eleman başına tekdüze mi? Bir aşama milyonlarca piksele aynı aritmetiği uyguluyorsa GPU biçimlidir. Her iterasyon farklı dallanıyorsa, her biri kendi durum makinesini çalıştıran onlarca takip gibi, CPU biçimlidir ve OpenMP’ye gider.
İkincisi: paralellik gerçekte ne kadar geniş? GPU’nun başa baş noktası binlerce bağımsız eleman ister. Piksel başına iş bu sınavı geçer. Takip başına iş geçmez: yoğun bir sahne sana onlarca nesne verir ve hiçbir GPU bunu donanımına yayamaz. Bir avuç CPU çekirdeği o işe fazlasıyla yeter.
Üçüncüsü, ki insanların atladığı bu: veri aşamadan önce ve sonra nerede yaşıyor? GPU’da hızlı hesaplayan ama hattın ortasında kare boyutunda bir indirme dayatan aşama net zarardır. Sırf kareyi zaten olduğu yerde tuttuğu için işi birden fazla kez “yanlış” işlemciye taşıdım. Taşıma faturası aşamanın maliyetinin parçasıdır ve çoğu zaman hesabın kendisinden ağır basar.
Vaka: savunma hattını aşama aşama yönlendirmek
Bütün bunların teori olmaktan çıktığı yer savunma takip projesi: görünür ve termal kanallar üzerinde gerçek zamanlı tespit ve takip, CUDA C ve OpenMP ile C++, Jetson üzerinde, operatörün karşısında bir Qt arayüzü. Oradan sayı, gecikme ya da test verisi paylaşamam. Yönlendirme mantığını paylaşabilirim; çünkü herhangi bir hatta taşınabilen kısım o.
Yakalama ve çözme CPU’da kalıyor. Sensörler kareyi nasıl veriyorsa öyle alırsın; o kod G/Ç ve üretici API çağrısından ibaret, paralelleştirilecek bir şey yok. Oradan kare GPU’ya bir kez çıkıyor ve ön işleme indiği yerde oluyor: distorsiyon düzeltme, yeniden boyutlandırma, biçim dönüşümü, ağ için normalizasyon. Bunların her biri, hem görünür hem termal kanalda, her piksele tekrarlanan aynı aritmetik. Termal kare tek kanallı ve kendi koşullamasını istiyor; ama işin biçimi aynı, o yüzden aynı CUDA yolunda gidiyor.
Çıkarım da GPU’da, doğrudan o tamponlardan çalışıyor. Sonra bütün tasarımın dönüm noktası geliyor: GPU’dan geri dönen şey kutular, skorlar ve sınıflar. Megabayt değil kilobayt. O çizginin aşağısındaki her şey nesne başına mantık: tespitlerle mevcut takipler arasında eşleştirme, filtre güncellemeleri, takip doğumu ve ölümü, kapılama kararları. Onlarca eleman üzerinde dallanma dolu kod; bir warp için tam yanlış, CPU çekirdeklerine yayılmış OpenMP için tam doğru biçim.
Kare: projenin gerçek çıktısı.
Kopya sorusu kararları sürekli şekillendirdi. Kalıp şöyleydi: hattın ortasında tam çözünürlüklü kareyi CPU’ya geri isteyen her şeye tasarım kokusu muamelesi yaptım ve yalnızca gösterim yolu tam kareye dokunana kadar yapıyı değiştirdim. Operatör katmanı, kutular, nişangâh, semboller, zaten ekrana giden videonun üzerine çiziliyor; yani fazladan kare trafiği eklemiyor. Qt arayüzü de tek kurallı kendi iş parçacığında yaşıyor: arayüz hattı asla bloklamaz, hat da arayüzü. OpenMP paralel bölümleri yakalama, işleme ve gösterimi üst üste bindiriyor; böylece her aşama komşusunun gecikmesini gizliyor.
Bu disiplinin sonucu buraya basabileceğim bir benchmark değil. Sonuç şu: aynı yönlendirme, mimarinin değişmesine gerek kalmadan yıllarca saha kullanımına dayandı. Pikseller GPU’ya, nesneler CPU’ya; aralarında da yalnızca küçük şeylerin geçmesine izin verilen sıkı bir sınır.
Saat tek, o yüzden ölç
Amdahl yasası hepsinin altındaki sıkıcı gerçek: paralel kısmı istediğin kadar hızlandır, tavanını seri tutkal belirler. Pratikte darboğazım neredeyse hiç gurur duyduğum çekirdek olmadı. Ya unuttuğum bir kopyaydı ya da GPU’yu besleyen iş parçacığıyla çekirdek kavgası eden bir OpenMP bölgesi.
İki alışkanlık işe yarıyor. OpenMP’nin gerçekte kaç iş parçacığı aldığını sabitle; her şeyi kapıp sürecin geri kalanını aç bırakmasın. Ve hiçbir şeyi optimize etmeden önce profil çıkar: Jetson’da tegrastats ve her aşamanın etrafına konmuş bir zamanlayıcı, neyin “hızlı olması gerektiğine” dair her sezgiden daha çok şey söyler.
Yani cevap OpenMP mi CUDA mı değil. OpenMP ve CUDA; her biri kendi biçimindeki paraleli yapıyor ve ikisinin faturası da aynı kare süresine kesiliyor.