“CUDA kullanıyoruz” cümlesini mülakatlarda, kod incelemelerinde ve tedarikçi toplantılarında duyuyorum ve üç bambaşka anlama gelebiliyor. Bazen PyTorch’un GPU’da koştuğu anlamına geliyor. Bazen birinin OpenCV CUDA modülüne bağlandığı anlamına geliyor. Bazen elle yazılmış çekirdekler demek. Bunlar maliyetleri farklı, ayrı işler; savunma tarafında, Jetson üzerinde gerçek zamanlı tespit ve takip için C++ hatları yazdığım yerde, yanlış katmanı seçmek haftaların nasıl kaybedildiğidir.
Şekil: kendi çizimim.
İşte birinin bana zamanında çizmesini istediğim harita.
CUDA bir platformdur, dil değil
CUDA, NVIDIA’nın gönderdiği yığının tamamı: sürücü, çalışma zamanı, nvcc derleyicisi ve üstüne cuBLAS, cuDNN, cuFFT, NPP gibi bir yığın kütüphane. TensorRT de o yığının üstünde oturuyor. Bir YOLO modelini GPU’da koşturduğunda, tek satır yazmadan bu anlamda CUDA kullanıyorsun. “CUDA kullanan” insanların çoğu tamamen bu katmanda yaşıyor ve bunda sorun yok. Çoğu iş için doğru katman bu.
CUDA C: çekirdekleri sen yazarsın
Bir katman aşağıda CUDA C var: __global__ fonksiyonlar yazdığın, grid, blok
ve thread cinsinden düşündüğün C++ uzantısı. Tam kontrol ve tam sorumluluk
alırsın: bellek erişim düzeni, doluluk, paylaşımlı bellek, sadece yük altında
ortaya çıkan yarış durumları.
Ne zaman değer? Benim tecrübemde iki durumda. Birincisi, hiçbir kütüphanenin sunmadığı bir işlem. İkincisi, bellek gidiş dönüşlerini öldürmek için adımları kaynaştırmak. Bizim ön işleme zinciri, yeniden boyutlandırma artı renk dönüşümü artı normalizasyon, üç kütüphane çağrısı olarak bellekten üç kez geçiyordu. Tek çekirdeğe kaynaştırınca tek geçiş oldu. Bu tür kazanç gerçek, ama ona profiler söylediği için uzanırsın, çekirdek yazmak daha ciddi hissettirdiği için değil.
OpenCV CUDA modülü: başkasının çekirdekleri
Arada cv::cuda duruyor. Kareyi bir GpuMat içine yüklersin, CPU API’sini
aynalayan fonksiyonları çağırırsın, sonucu indirirsin. Klasik işlemlerde
kapsam fena değil: filtreleme, bükme, optik akış, arka plan çıkarma.
Sahadan iki uyarı. Bir: pip paketleri bu modülü içermiyor. OpenCV’yi CUDA açık şekilde kendin derliyorsun ve Jetson’da o derleme başlı başına bir öğleden sonra. Kart kurulumunu JetPack yazısında anlattım. İki: aynalanmış API, aynalanmış performans demek değil. Bazı fonksiyonlar CPU’dan dramatik hızlı, bazıları ucu ucuna hızlı, birkaçı küçük görüntülerde daha yavaş. İmanla değil, kendi kare boyutunda ölçerek karar ver.
Asıl maliyet hesap değil, transfer
Üç katmanı birleştiren hata: GPU’ya arada bir çağırdığın hızlı bir fonksiyon gibi davranmak. Her yükleme ve indirme bir veri yolundan geçer ve iki kopya arasına sıkışmış hızlı bir çekirdek, yavaş bir hattır. Çözüm mimari. Yapabiliyorsan çözmeyi GPU’da yap, kareyi orada tut, bütün zinciri orada koştur ve sadece sonuçları indir; bir dedektör için bu bir avuç kutudur, görüntü değil.
Jetson bunu iyi yönde karmaşıklaştırıyor: CPU ve GPU fiziksel belleği paylaşıyor, yani sıfır kopya yollar mevcut. Ama otomatik değiller. Doğru ayırma tipini yine senin seçmen gerekiyor ve varsayılan nadiren odur.
Ben gerçekte nasıl seçiyorum
Sıram sıkıcı ve işe yarıyor. Önce kütüphane: OpenCV CUDA, NPP, çıkarım için TensorRT. Özel CUDA C sadece profiler’ın gösterdiği yerde; bu şimdiye kadar ön işleme kaynaştırması ve bir son işleme adımı demek oldu. Ve dürüst olalım, bazen cevap CPU: küçük görüntülerde ya da mantık ağırlıklı kodda birkaç OpenMP thread’i, transfer maliyetini toptan yener. O dengeyi OpenMP ile CUDA farkı yazısında anlattım.
Hangi katmanda durduğunu bil. Gördüğüm GPU “optimizasyon” başarısızlıklarının çoğu yavaş çekirdek değildi. Yanlış katmandaki doğru koddu.