6. Paralel Çalışma ve Doğruluk

---

§15 Execution

WGSL modülündeki bir shader, GPU üzerinde çok sayıda invocation (çağrı) olarak çalıştırılır. Bu bölüm, invocation'ların bireysel ve kolektif olarak nasıl yürütüldüğünü tanımlar.

---

15.1 Program Order Within an Invocation

Her bir invocation içinde, WGSL statement'ları kontrol akışı sırasına (control flow order) göre yürütülür.

• Bir statement, yürütme sırasında sıfır veya daha fazla kez çalıştırılabilir.
• Her bir yürütme, benzersiz bir dynamic statement instance oluşturur.
• İfade (expression) değerlendirme sırası soldan sağa (left-to-right):

// foo() her zaman bar()'dan ÖNCE değerlendirilir
let result = foo() + bar();

• İç içe geçmiş ifadeler veri bağımlılıkları (data dependencies) oluşturur — iç ifade, dış ifadeden önce değerlendirilmelidir.

---

15.2 Uniformity

Uniformity (tekdüzelik), GPU programlamanın en kritik kavramlarından biridir. Belirli operasyonlar — collective operation'lar (barrier, derivative, texture sampling) — tüm invocation'ların eş zamanlı olarak çalıştırmasını gerektirir.

Temel Kavramlar

• Uniform control flow: Aynı kapsamdaki tüm invocation'lar, programın aynı noktasını eş zamanlı olarak yürütür.
• Non-uniform control flow: İnvocation'ların bir alt kümesi operasyonu yürütürken diğerleri yürütmez. Bu durum, yanlış veya taşınabilir olmayan davranışa yol açar.

Non-uniform kaynakları

Uniformity Failure

WGSL derleyicisi, her collective operasyonun uniform control flow içinde çalışıp çalışmadığını doğrulamak için statik uniformity analizi yapar. Doğrulama başarısız olursa:

15.2.1 Terminology and Concepts

Belirli bir invocation grubu için:

• Uniform control flow: Verilen uniformity scope (tekdüzelik kapsamı) içindeki tüm invocation'lar, programın aynı noktasını sanki adım adım (lockstep) yürütüyormuş gibi çalışır.
• Uniformity scope'lar:

• Uniform value: Uniform control flow'da çalıştırılan ve tüm invocation'ların aynı değeri hesapladığı bir ifade.
• Uniform variable: Her canlı olduğu noktada tüm invocation'ların aynı değeri tuttuğu bir yerel değişken.

15.2.2 Uniformity Analysis Overview

Uniformity analizi, her fonksiyonda iki şeyi doğrular:

1. Fonksiyon, çağırdığı diğer fonksiyonlar için uniformity gereksinimlerini karşılar. 2. Fonksiyon, her çağrıldığında uniformity gereksinimlerini karşılar.

Analiz özellikleri:

• Sound (sağlam): Uniformity gereksinimlerini ihlal eden bir program için her zaman hata tetiklenir.
• Refactoring-safe: Kodu fonksiyona çıkarmak veya fonksiyonu inline yapmak, geçerli bir shader'ı geçersiz kılmaz.
• İzlenebilir: Hata durumunda, root cause'a kadar uzanan bir implikasyon zinciri oluşturulabilir.

Analiz, çağrı grafiğinde yaprak fonksiyonlardan entry point'e doğru (topolojik sırada) ilerler. WGSL'de recursion yasak olduğu için döngü riski yoktur.

Potential-Trigger-Set

Her fonksiyon çağrısı için, çağrı non-uniform control flow'da yapılırsa tetiklenecek olan triggering rule'lar hesaplanır:

15.2.3 Analyzing the Uniformity Requirements of a Function

Her fonksiyon iki aşamada analiz edilir:

1. Graf oluşturma: Fonksiyonun sözdizimi üzerinden yönlü bir graf kurulur. 2. Graf keşfi: Bu graf üzerinden kısıtlamalar hesaplanır ve olası uniformity failure tespit edilir.

Grafın Özel Düğümleri

Temel kural: RequiredToBeUniform.error → MayBeNonUniform arasında bir yol varsa, uniformity violation tetiklenir.

Fonksiyon Tag'leri

Her fonksiyon için hesaplanan tag'ler:

15.2.4 Pointer Desugaring

Uniformity analizi öncesinde, function adres uzayındaki pointer parametreler yerel değişken bildirimi olarak çözümlenir (desugar):

// Orijinal
fn foo(p: ptr<function, array<f32, 4>>, i: i32) -> f32 {
  let p1 = p;
  var x = i;
  let p2 = &((*p1)[x]);
  x = 0;
  *p2 = 5;
  return (*p1)[x];
}

// Analiz için eşdeğer form
fn foo_for_analysis(p: ptr<function, array<f32, 4>>, i: i32) -> f32 {
  var p_var = *p;             // p için yerel değişken
  let p1 = &p_var;            // p1 artık p_var'a referans
  var x = i;
  let x_tmp1 = x;             // x'in o anki değerini yakala
  let p2 = &(p_var[x_tmp1]);  // p1'in initializer'ını substitüe et
  x = 0;
  *(&(p_var[x_tmp1])) = 5;    // p2'nin initializer'ını substitüe et
  return (*(&p_var))[x];      // p1'in initializer'ını substitüe et
}

Bu dönüşüm, pointer'ın root identifier'ını her kullanım noktasında doğrudan ortaya çıkararak analizi basitleştirir.

15.2.5 Function-scope Variable Value Analysis

Fonksiyon kapsamındaki değişkenlerin değeri, ona ulaşan atamalar (assignments) üzerinden analiz edilir.

Full/Partial reference kavramı:

• Full reference: Bir değişkenin tüm bellek konumlarına erişim (v ifadesi)
• Partial reference: Yalnızca bir alt kümeye erişim (v.member, arr[0])

15.2.6 Uniformity Rules for Statements

Statement'lar için uniformity kuralları, gelen kontrol akışı düğümünü (CF) alır ve çıkan kontrol akışı düğümünü üretir. Temel kurallar:

15.2.7 Uniformity Rules for Function Calls

Fonksiyon çağrıları, çağrılan fonksiyonun tag'lerine göre analiz edilir:

• CallSiteRequiredToBeUniform.S → Çağrı noktasının kontrol akışı RequiredToBeUniform.S'e bağlanır
• ParameterRequiredToBeUniform.S → Argümanın değer düğümü RequiredToBeUniform.S'e bağlanır
• ReturnValueMayBeNonUniform → Sonuç düğümü MayBeNonUniform'a bağlanır
• ParameterReturnContentsRequiredToBeUniform → Parametre, dönüş değerinin uniformity'sine etki eder

15.2.8 Uniformity Rules for Expressions

İfade analizi, hem kontrol akışı (CF) hem de değer (V) düğümü üretir:

15.2.9 Annotating the Uniformity of Every Point

Analiz tamamlandığında, programın her noktasındaki uniformity durumu graftan çıkarılabilir:

• RequiredToBeUniform.S → MayBeNonUniform yolu: uniformity failure
• Yol üzerindeki düğümler: hata mesajındaki açıklama zinciri

15.2.10 Examples

15.2.10.1 Invalid textureSample Function Call

@group(0) @binding(0) var t: texture_2d<f32>;
@group(0) @binding(1) var s: sampler;

@fragment
fn main(@builtin(position) pos: vec4<f32>) {
  if pos.x > 1.0 {
    // ❌ GEÇERSİZ: textureSample derivative hesaplar;
    // pos non-uniform → if koşulu non-uniform → non-uniform control flow
    _ = textureSample(t, s, pos.xy);
  }
}

Neden geçersiz? position built-in değeri non-uniform → pos.x > 1.0 koşulu non-uniform → if dalları non-uniform control flow oluşturur → textureSample (derivative hesaplar) non-uniform control flow'da çağrılır.

15.2.10.2 Function-scope Variable Uniformity

@group(0) @binding(0) var<storage, read_write> a: i32;
@group(0) @binding(1) var<uniform> b: i32;

@compute @workgroup_size(16, 1, 1)
fn main() {
  var x: i32;
  x = a;
  if x > 0 {
    // ❌ GEÇERSİZ: x, mutable modül-kapsamı 'a'dan türetildi
    workgroupBarrier();
  }
  x = b;
  if x < 0 {
    // ✅ GEÇERLİ: x artık immutable uniform 'b'den türetildi
    storageBarrier();
  }
}

Bu örnek, value analysis'in bir değişkenin ömrü boyunca farklı uniformity dönemlerini ayırt edebildiğini gösterir. İlk if'ten sonra x yeniden atandığı için, ikinci if bağımsız olarak analiz edilir.

15.2.10.3 Composite Value Analysis Limitations

struct Inputs {
  @builtin(workgroup_id) wgid: vec3<u32>,           // Uniform
  @builtin(local_invocation_index) lid: u32   // Non-uniform
}

@compute @workgroup_size(16, 1, 1)
fn main(inputs: Inputs) {
  // ❌ GEÇERSİZ: inputs composite olarak non-uniform olarak işaretlenir
  // (lid non-uniform olduğu için tüm inputs non-uniform)
  if inputs.wgid.x == 1 {
    workgroupBarrier();
  }
}

Çözüm: Composite'ı ayırın:

@compute @workgroup_size(16, 1, 1)
fn main(@builtin(workgroup_id) wgid: vec3<u32>,
        @builtin(local_invocation_index) lid: u32) {
  // ✅ GEÇERLİ: wgid ayrı parametre → uniform olarak izlenir
  if wgid.x == 1 {
    workgroupBarrier();
  }
}

15.2.10.4 Uniformity in a Loop

@compute @workgroup_size(16, 1, 1)
fn main(@builtin(local_invocation_index) lid: u32) {
  for (var i = 0u; i < 10; i++) {
    workgroupBarrier();  // ❌ GEÇERSİZ
    if (lid + i) > 7 {
      break;  // Non-uniform break → sonraki iterasyonlarda CF non-uniform
    }
  }
}

Neden? lid non-uniform → break koşulu non-uniform → bazı invocation'lar döngüden erken çıkar → sonraki iterasyonlarda workgroupBarrier non-uniform control flow'da çağrılır.

15.2.10.5 User-defined Function Calls

fn scale(in1: f32, in2: f32) -> f32 {
  let v = in1 / in2;
  return v;
}

@group(0) @binding(0) var t: texture_2d<f32>;
@group(0) @binding(1) var s: sampler;

@fragment
fn main(@builtin(position) pos: vec4<f32>) {
  let tmp = scale(pos.x, 0.5);
  if tmp > 1.0 {
    // ❌ GEÇERSİZ: scale'in dönüş değeri her iki parametreye bağlıdır
    // (ParameterReturnContentsRequiredToBeUniform);
    // pos.x non-uniform → tmp non-uniform → non-uniform control flow
    _ = textureSample(t, s, pos.xy);
  }
}

Analiz, scale fonksiyonunun her iki parametresini ParameterReturnContentsRequiredToBeUniform olarak etiketler — dönüş değerinin uniform olması için parametrelerin de uniform olması gerekir.

---

15.3 Compute Shaders and Workgroups

Workgroup Kavramı

Workgroup, bir compute shader entry point'ini eş zamanlı olarak çalıştıran ve workgroup adres uzayındaki değişkenleri paylaşan bir invocation kümesidir.

Workgroup Grid

Compute shader'ın workgroup grid'i, @workgroup_size(x, y, z) attribute'u ile tanımlanan 3D tamsayı koordinat uzayıdır:

• 0 ≤ i < workgroup_size_x
• 0 ≤ j < workgroup_size_y
• 0 ≤ k < workgroup_size_z

Her grid noktası için tam olarak bir invocation bulunur.

Invocation ID'leri

Dispatch

Dispatch komutu bir dispatch size (group_count_x, group_count_y, group_count_z) belirtir. Compute shader grid:

• 0 ≤ CSi < workgroup_size_x × group_count_x
• 0 ≤ CSj < workgroup_size_y × group_count_y
• 0 ≤ CSk < workgroup_size_z × group_count_z

Garanti Verilmeyenler

WebGPU aşağıdaki garantileri vermez:

• ❌ Farklı workgroup'ların eş zamanlı çalışması
• ❌ Bir workgroup çalışırken diğerlerinin bloklanması
• ❌ Workgroup'ların belirli bir sırada başlatılması

---

15.4 Fragment Shaders and Helper Invocations

Quad Kavramı

Fragment shader invocation'ları, komşu position değerlerine sahip 2×2 grid'ler halinde organize edilir. Her grid bir quad olarak adlandırılır.

Helper Invocations

Bir grafik primitive'in kenarında, quad'ı doldurmak için yeterli RasterizationPoint olmayabilir. Bu durumda GPU, boş pozisyonlar için helper invocation'lar oluşturur.

Helper invocation kısıtlamaları:

Helper invocation'lar yalnızca derivative hesaplamaları için yardımcı olarak bulunur.

---

15.5 Subgroups

Subgroup, bir compute veya fragment shader'da eş zamanlı çalışan ve verimli veri paylaşımı yapabilen bir invocation kümesidir.

Subgroup Size

Subgroup ID'leri

Divergence

Aynı subgroup'taki invocation'lar farklı kontrol akışı yolları izlediğinde, subgroup yürütmesinin diverge olduğu söylenir:

• Active invocation: Bir subgroup operasyonunu eş zamanlı olarak yürüten invocation'lar
• Inactive invocation: Aynı operasyonu yürütmeyen invocation'lar (ya farklı dalda ya da subgroup boyutunu aşan "hayalet" invocation'lar)
• Helper invocation'lar cihaza bağlı olarak active veya inactive olabilir.

---

15.6 Collective Operations

15.6.1 Barriers

Barrier, bir programdaki bellek operasyonlarını sıralayan bir senkronizasyon fonksiyonudur.

Control barrier: Aynı workgroup'taki tüm invocation'lar tarafından eş zamanlı olarak yürütülür. Bu nedenle barrier'lar yalnızca uniform control flow içinde çağrılabilir.

@compute @workgroup_size(64)
fn main() {
  // Workgroup belleğine yazma
  shared_data[local_idx] = compute_value();

  // Tüm invocation'ların yazmasını bekle
  workgroupBarrier();

  // Artık diğer invocation'ların yazdığı değerler güvenle okunabilir
  let neighbor = shared_data[local_idx ^ 1];
}

WGSL'de üç barrier fonksiyonu vardır:

15.6.2 Derivatives

Partial derivative (kısmi türev), bir değerin bir eksen boyunca değişim oranıdır. Fragment shader'daki invocation'lar aynı quad içinde işbirliği yaparak yaklaşık türevler hesaplar.

Derivative hesaplayan fonksiyonlar

İmplicit derivative (fragment coordinate üzerinden):

• textureSample, textureSampleBias, textureSampleCompare

Explicit derivative (kullanıcı belirtilen değer üzerinden):

Uniformity gereksinimi: Derivative fonksiyonları uniform control flow içinde çağrılmalıdır. Aksi halde derivative_uniformity diagnostic'i tetiklenir ve sonuç indeterminate value olur.

15.6.3 Subgroup Operations

Subgroup built-in fonksiyonları, aynı subgroup'taki invocation'lar arasında verimli iletişim ve hesaplama sağlar. SIMT (Single Instruction Multiple Thread) operasyonlarıdır.

Active invocation'lar işbirliği yaparak sonuçları belirler. Taşınabilirlik, tüm invocation'ların active olduğu durumda (yani uniform control flow'da) maksimize edilir.

15.6.4 Quad Operations

Quad built-in fonksiyonları, bir quad (2×2 invocation grubu) üzerinde çalışır. Quad içi veri iletişimi için kullanışlıdır.

Subgroup operasyonlarındakine benzer şekilde, taşınabilirlik uniform control flow'da maksimize edilir.

---

15.7 Floating Point Evaluation

WGSL'nin floating point özellikleri IEEE-754 standardına dayanır; ancak GPU'ların yaptığı ödünleşimleri yansıtacak şekilde azaltılmış işlevsellik ve ek taşınabilirlik güvenceleri içerir.

15.7.1 Overview of IEEE-754

IEEE-754 binary floating point tipi, genişletilmiş gerçel sayı (extended real number) doğrusunu yaklaştırır:

Değer Kategorileri

Bit Temsili

Her IEEE-754 floating point değeri üç bitfield'den oluşur (MSB → LSB):

IEEE-754 Tipleri

Değer Yorumlama Algoritması

Bit temsilinden değer:

• Exponent tüm 1'ler → T=0 ise ±∞, T≠0 ise NaN
• Exponent tüm 0'lar → (−1)^Sign × 2^(−bias) × T × 2^(−tsw+1) → T=0 ise sıfır, T≠0 ise subnormal
• Diğer → (−1)^Sign × 2^(E−bias) × (1 + T × 2^(−tsw)) → normal

Temel Kavramlar

IEEE-754 İstisnalar

15.7.2 Differences from IEEE-754

WGSL, IEEE-754'ten şu şekillerde ayrılır:

Finite Math Assumption

WGSL'nin en önemli kurallarından biri:

Shader yürütmesi öncesinde (const/override expression'lar):

• Overflow, ∞ veya NaN → hata üretilir (shader-creation veya pipeline-creation error)

Shader yürütmesi sırasında (runtime expression'lar):

• Implementasyon overflow, ∞ ve NaN'ın olmadığını varsayabilir
• Bu durumda sonuç indeterminate value olur

Flush to Zero Kuralları

15.7.3 Floating Point Rounding and Overflow

Overflow durumunda sonuç, sonsuzluğa veya en yakın sonlu değere yuvarlanabilir. Bu, ara sonucun büyüklüğüne ve değerlendirmenin zamanlamasına bağlıdır.

Yuvarlama kuralları (X ara sonuç, tip T):

Son değer (X''):

• X' sonsuz veya NaN ise:
• const-expression → shader-creation error
• override-expression → pipeline-creation error
• runtime expression → indeterminate value
• Aksi halde: X'' = X'

15.7.4 Floating Point Accuracy

Correctly Rounded

Tam sonuç x için, correctly rounded sonuç:

• x tipin bir değeriyse → x
• Aksi halde → x'in üstündeki en küçük veya altındaki en büyük representable değer

ULP (Units in the Last Place)

ULP(x), x etrafındaki iki ardışık floating point değer arasındaki minimum mesafedir. Doğruluk sınırlarının ölçü birimi olarak kullanılır.

15.7.4.1 Accuracy of Concrete Floating Point Expressions

Temel aritmetik operatörler:

Seçili built-in fonksiyonların doğruluğu:

Inherited accuracy — bir operasyonun doğruluğu, alternatif bir WGSL ifadesinin doğruluğundan devralınır. Örneğin tan(x) → sin(x) / cos(x).

15.7.4.2 Accuracy of AbstractFloat Expressions

AbstractFloat operasyonlarının doğruluğu:

15.7.5 Reassociation and Fusion

Reassociation (Yeniden İlişkilendirme)

Operasyonların matematiksel olarak eşdeğer, ancak farklı sırada hesaplanması:

(a + b) + c  →  a + (b + c)    // Reassociation
(a - b) + c  →  (a + c) - b    // Reassociation
(a * b) / c  →  (a / c) * b    // Reassociation

Fusion (Birleştirme)

Birden fazla operasyonun tek bir daha doğru operasyona birleştirilmesi:

• Birleştirilen ifade, orijinalinden en az eşit doğruluğa sahip olmalıdır
• Örnek: FMA (fused multiply-add) ayrı çarpma ve toplamadan daha doğru olabilir

15.7.6 Floating Point Conversion

Float → Integer

X (float) → T (integer tipi)

Örnekler:

• 3.9f → u32 = 3u
• -1f → u32 = 0u
• 1e20f → u32 = 4294967040u (en büyük f32-representable u32 değeri)
• -3.9f → i32 = -3i

Integer/Float → Float

15.7.7 Domains of Floating Point Expressions and Built-in Functions

Her floating point operasyonun bir domain'i (tanım kümesi) vardır:

• Kısıtlama belirtilmemişse: domain total (tüm sonlu ve sonsuz girdiler)
• Aksi halde: açık domain listelenir (ör: acos domain'i [−1, 1])

Component-wise operasyonlar için yalnızca skaler durum tanımlanır; vektör durumu component-wise semantiklerden türetilir.

Inherited accuracy olan operasyonlarda domain:

• Açık olarak tanımlanır, veya
• Implied from linear terms — "inherited from" ifadesindeki toplama, çıkarma ve çarpma operasyonlarının domain kısıtlamalarının birleştirilmesiyle türetilir.

Örnek: dot(a, b) (2-element) = a[0]*b[0] + a[1]*b[1]:

• Çarpma: 0 × ∞ ve ∞ × 0 hariç
• Toplama: Zıt işaretli sonsuzların toplamı hariç
• Sonuç domain: Bu dışlamaların birleşimi

---