CUDA'da thread, block ve grid neyi anlatıyor?

Bir CUDA kernel’ini okumaya başlarken ilk soru şu: Bu thread hangi veriden sorumlu? Thread, block ve grid kavramları bu eşlemeyi kurmak için kullanılıyor.

Bu yazı, NVIDIA belgelerine dayanan bir kavram anlatımı. Buradaki sayısal örnek bir indeks hesabı; donanım üzerinde elde edilmiş bir performans sonucu değil.

Kernel tek bir işi nasıl çoğaltıyor?

CPU’da çalışan host kodu, GPU’da çalışacak bir fonksiyonu başlatır. Bu fonksiyona kernel denir. Aynı kernel kodu birçok thread için çalışır; her thread kendi indekslerini kullanarak verinin ilgili bölümünü işler.

Thread’ler block grupları içinde yer alır. Bir kernel çağrısındaki block’ların bütünü grid oluşturur. Bir boyutlu bir örnekte şu üç değeri kullanırız:

Değer	Anlamı
`threadIdx.x`	Thread’in kendi block’u içindeki indeksi.
`blockIdx.x`	Block’un grid içindeki indeksi.
`blockDim.x`	Bir block’taki thread sayısı.

Bu düzen, veriyi thread’lere bölmek için bir adresleme sistemi sağlar. Block sayısı ile GPU’daki fiziksel işlem birimi sayısı aynı kavram değildir. Modelin ayrıntıları NVIDIA Programming Guide’da açıklanıyor.

Örnek: her thread bir elemanı toplasın

Aşağıdaki çekirdek parçasında c[i] = a[i] + b[i] hesabını thread’lere dağıtıyoruz. a, b ve c adreslerinin GPU tarafından erişilebilir olması, n elemanlık alanı göstermesi ve çıktı alanının eşzamanlı erişim çatışması yaratmaması gerekiyor.

#include <cstddef>

__global__ void add(const float* a, const float* b,
                    float* c, std::size_t n) {
    const std::size_t i =
        static_cast<std::size_t>(blockIdx.x) * blockDim.x
        + threadIdx.x;

    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

Bu kod yalnızca kernel tanımını gösteriyor; bellek ayırma, veri kopyalama, çağrı ve hata kontrolünü içeren tam bir program değil.

İndeks hesabını küçük bir örnekle izleyelim: n = 1000 için block başına 256 thread seçersek dört block, toplam 1.024 thread oluşturur. İndeksi 1.000-1.023 arasında olan 24 thread, if (i < n) nedeniyle diziye erişmez.

Örneğin üçüncü block’taki, yani blockIdx.x = 2 olan block’taki 7 numaralı thread, 2 * 256 + 7 = 519 elemanını işler. Burada 256 bir örnek seçim; her kernel için en iyi block boyutu olduğu iddiası yok.

Komşu thread, komşu adres

Bu örnekte ardışık thread’ler ardışık float adreslerine erişir. Bir warp içindeki bellek erişimlerinin az sayıda işlemde karşılanabilmesine coalescing denir. Hizalama ve gerçek erişim örüntüsü sonucu etkiler; “ardışık” demek her koşulda tek bellek işlemi demek değildir. NVIDIA Best Practices Guide bu ayrımı örneklerle gösteriyor.

Çağrı süresi ile işin bitmesi aynı şey mi?

Kernel başlatmak genellikle host açısından asenkrondur. CPU’daki bir kronometreyi yalnızca başlatma çağrısının etrafına koymak, GPU işinin tamamını ölçmeyebilir.

GPU zamanlaması için aynı stream’e kaydedilen CUDA event’leri kullanılabilir; bitiş event’inin tamamlanması beklenir ve aradaki süre alınır. Verinin CPU’dan GPU’ya gidişi ve geri dönüşü bu sınırın dışında bırakılmışsa, ölçüm uçtan uca süreyi göstermez. Zamanlama rehberi bu ayrıntıları açıklıyor.

Bir kernel’i okurken önce indeksin doğru elemana gittiğini, sonra erişim düzenini ve en son zamanlanan bölgenin ne olduğunu sormak iyi bir başlangıç. Performans iddiası için bunlara gerçek girdi, cihaz, derleme ayarları ve tekrar kayıtları da eklenmeli.

END OF NOTE / KEEP EXPLORING

Small experiments. Honest notes.