Bölüm 07

Atomik ve Bellek Senkronizasyonu

Atomic işlemler, çok sayıda invocation aynı bellek adresine dokunduğunda gerekir. Barrier'lar ise bir workgroup içinde net bir önce/sonra sınırı koyar. Bu bölüm, GPU kodunun bağımsız lane'lerden paylaşılan state'e geçtiği noktayı öğretir.

Atomic tipler

atomic<T>, aynı değere birden fazla invocation erişebildiğinde anlam kazanır. WGSL tarafında bileşen tipi i32 veya u32 olur; atomic değerler pratikte workgroup veya storage address space'lerinde karşımıza çıkar.

var<workgroup> local_count : atomic<u32>;

@group(0) @binding(0)
var<storage, read_write> global_count : atomic<u32>;

workgroup atomic, aynı workgroup içindeki lane'ler işbirliği yapıyorsa uygundur. storage atomic ise birçok workgroup'un host tarafından görülebilen ortak bir buffer'a birikim yaptığı durumlarda kullanılır.

Atomic fonksiyonlar

Her atomic fonksiyon tek bellek konumu için bölünmez bir güncelleme sağlar. En önemli alışkanlık şudur: hangi fonksiyon eski değeri döndürür, hangisi yalnızca yeni değeri yayımlar?

Klasik ilk örnek: histogram

Her invocation farklı input okuyabilir ama aynı bin'e denk gelebilir. Atomic yoksa iki lane aynı eski bin değerini okuyup artışlardan birini kaybettirebilir. atomicAdd, increment işlemini tek bölünmez güncelleme haline getirir.

@group(0) @binding(0)
var<storage, read> input : array<u32>;

@group(0) @binding(1)
var<storage, read_write> bins : array<atomic<u32>>;

@compute @workgroup_size(256)
fn histogram(@builtin(global_invocation_id) gid : vec3u) {
  let i = gid.x;
  if (i >= arrayLength(&input)) {
    return;
  }

  let bin = input[i] & 255u;
  atomicAdd(&bins[bin], 1u);
}

Barrier workgroup sınırı çizer

Barrier, aynı workgroup içindeki bütün invocation'ların verdiği bir sözdür: herkes bu noktaya gelmeden kimse ilerlemez. Bu yüzden barrier çağrıları uniform şekilde erişilebilir olmalıdır. Barrier tüm GPU işlerini bitiren genel bir düğme değildir; belirli scope içindeki lane'leri senkronize eder.

Workgroup scratch memory

En yaygın barrier deseni şudur: veriyi workgroup memory'ye al, bekle, reduce veya tekrar kullan, gerekiyorsa bir sonraki aşama için tekrar bekle. Her invocation algoritmanın bir lane'ini taşır ama array workgroup içinde paylaşılır.

var<workgroup> scratch : array<u32, 256>;

@group(0) @binding(0)
var<storage, read> input : array<u32>;

@group(0) @binding(1)
var<storage, read_write> partials : array<u32>;

@compute @workgroup_size(256)
fn sum_block(
  @builtin(local_invocation_id) lid : vec3u,
  @builtin(global_invocation_id) gid : vec3u,
  @builtin(workgroup_id) wid : vec3u
) {
  let local_index = lid.x;
  scratch[local_index] = select(0u, input[gid.x], gid.x < arrayLength(&input));
  workgroupBarrier();

  var stride = 128u;
  loop {
    if (local_index < stride) {
      scratch[local_index] = scratch[local_index] + scratch[local_index + stride];
    }
    workgroupBarrier();

    if (stride == 1u) {
      break;
    }
    stride = stride / 2u;
  }

  if (local_index == 0u) {
    partials[wid.x] = scratch[0];
  }
}

Memory layout hala önemlidir

Atomic ve barrier senkronizasyon problemini çözer; layout kurallarını ortadan kaldırmaz. Buffer içinde struct varsa host ve shader aynı alignment/stride hikayesinde buluşmalıdır.

Çalışma döngüsü

1. Paylaşılan adresi adlandır: workgroup memory mi, storage memory mi, ikisi birden mi?
2. Bütün yazmaları bul; iki invocation aynı adrese denk gelebilir mi karar ver.
3. Adres çakışabiliyorsa update'e uyan en küçük atomic operasyonu seç.
4. Sonraki kod workgroup memory okuyorsa, bütün lane'lerin ulaşacağı barrier noktasını koy.
5. Studio'da legality için diagnostics, API sözleşmesi için reflection, sonuç için readback kullan.