Warp Divergence

GPU’da thread’ler warp adı verilen 32 thread‘lik gruplar halinde yürütülür. Bir warp içindeki thread’ler aynı instruction stream’i takip ederek farklı data üzerinde çalışır.

Buraya kadar her şey güzel. Sıkıntı, farklı instructionlar çalışması gerektiğinde başlıyor. Örneğin if() koşulu koyduğumuzda threadlerin bir kısmı if bloğunu bir kısmı else bloğunu çalıştıracak. Fakat bir warp içinde her iki bloğun instruction’larını aynı anda çalıştıramayız.

Benzetme yapalım

Bir komutanız ve emrimizde 32 adet asker var. Kırmızı bayrak tutan askerleri 2 adım öne beyaz bayrak tutan askerleri 3 adım geriye almak istiyoruz. Bu durumda komutlarımız şöyle olurdu

1. Kırmızı bayrak tutanlar 1 adım ileri marş
2. Kırmızı bayrak tutanlar 1 adım ileri marş
3. Beyaz bayrak tutanlar geriye dön
4. Beyaz bayrak tutanlar 3 adım ileri marş
5. Beyaz bayrak tutanlar geriye dön

• Beyaz bayrak tutan askerler, kırmızı hareket ederken hareketsizler
• Kırmızı bayrak tutanlar ise beyazlar hareket ederken hareketsiz kalıyorlar.
• Yani burada bazı askerler diğer askerleri beklemek zorunda kalıyor.
• YANİ BAZI CORE’LAR ÇALIŞIRKEN ÖTEKİLER BEKLİYOR

• Peki 2 grup askerimiz olsaydı? Yani 2 warp asker
• Her grup için 1 komutana ihtiyacımız olacaktı
• Artık işlem gücümüz daha fazla ama sıkıntımız aynı. Halen bazı askerler bazı komutlarda boşta bekliyor. Biz bunu istemiyoruz
• Peki bu durumda ne yapabiliriz?

• Soldaki komutan beyaz bayraklı grubu,
• Sağdaki komutan kırmızı bayraklı grubu kontrol ediyor.

Soldaki komutan

1. 1 adım ileri marş
2. 1 adım ileri marş

Sağdaki komutan

1. Geriye dön
2. 3 adım ileri marş
3. Geriye dön

• Görüldüğü üzere askerleri düzgün bir şekilde gruplandırdığımızda boşta bekleme durumu gerçekleşmiyor. Aynı şekilde threadleri warp’larda doğru şekilde gruplandırdığımızda core’lar boşta kalmıyor.

Ne yapılabilir:

1. Branch’i warp sınırına hizala — Eğer koşul warp boyunca (32’nin katları) sabitse divergence olmaz. Örneğin threadIdx.x / 32 bazlı bir koşul, tüm warp aynı sonucu üretir.
2. Predication kullan — Basit koşullu işlemler için branch yerine aritmetik/mask kullan:

// Divergent:
if (cond) x = a; else x = b;
// Predicated:
x = cond * a + (1 - cond) * b;
// ya da
x = cond ? a : b; // compiler genelde bunu zaten predicate'e çevirir basit case'lerde

// Divergent:
if (cond) x = a; else x = b;
// Predicated:
x = cond * a + (1 - cond) * b;
// ya da
x = cond ? a : b; // compiler genelde bunu zaten predicate'e çevirir basit case'lerde

3. Data’yı yeniden düzenle (sort/reorder) — Eğer divergence data-dependent ise (örneğin bir threshold’a göre farklı işlem), veriyi önceden sınıflandırıp aynı path’e girecek thread’leri aynı warp’a toplayabilirsin (stream compaction, bucket sort).
4. Branch’i warp dışına taşı — Eğer mümkünse koşulu CPU tarafında veya kernel launch parametrelerinde çöz, farklı kernel’ler çağır (örneğin farklı case’ler için ayrı kernel launch).
5. Loop’larda early-exit’ten kaçın — Bir thread erken çıkarken diğerleri devam ediyorsa warp yine de tam süre bekler; loop trip count’u warp içinde eşitlemeye çalış.
6. __syncwarp() ve warp-level primitive’ler — Divergence sonrası gereksiz sync noktalarını azaltmak için warp-level intrinsic’leri (__shfl_sync, __ballot_sync) kullan, bunlar zaten divergence-aware.