Паралельні обчислення на GPU — SIMT, розгалуження варпів і модель Roofline

1. Виконання SIMD проти SIMT

2. Варпи, блоки потоків та сітки

3. Ієрархія пам’яті GPU

Конфлікти банків спільної пам’яті

4. Модель Roofline

Модель Roofline (Williams, Waterman, Patterson, 2009) — це простий візуальний інструмент, що допомагає зрозуміти, чи ядро обмежене обчисленнями, чи пам’яттю. Два «дахи» визначають стелю продуктивності:

5. Заповненість та обмежувальні ресурси

6. Обчислювальні шейдери WebGPU

WebGPU (доступний у Chrome 113+, Firefox Nightly) надає обчислювальні шейдери GPU у браузері за допомогою WGSL (WebGPU Shading Language):

// Обчислювальний шейдер WGSL — паралельна префіксна сума (один прохід)
@group(0) @binding(0) var<storage, read>  input  : array<f32>;
@group(0) @binding(1) var<storage, read_write> output : array<f32>;
@group(0) @binding(2) var<uniform> params : struct { n: u32 };

var<workgroup> shared_data : array<f32, 256>;

@compute @workgroup_size(256, 1, 1)
fn main(
  @builtin(global_invocation_id) gid : vec3<u32>,
  @builtin(local_invocation_id)  lid : vec3<u32>
) {
  let i = gid.x;
  let l = lid.x;

  // Завантаження у спільну пам’ять робочої групи
  shared_data[l] = select(0.0f, input[i], i < params.n);
  workgroupBarrier();

  // Паралельне сканування Hillis-Steele (log₂ 256 = 8 ітерацій)
  for (var offset = 1u; offset < 256u; offset <<= 1u) {
    if (l >= offset) {
      shared_data[l] += shared_data[l - offset];
    }
    workgroupBarrier();
  }

  if (i < params.n) { output[i] = shared_data[l]; }
}

Запускайте цей шейдер через passEncoder.dispatchWorkgroups(Math.ceil(n/256), 1, 1). Кожна робоча група обробляє 256 елементів за ~8 розділених бар’єрами проходів, причому спільна пам’ять усуває звернення до глобальної пам’яті між проходами.

7. Ключові патерни оптимізації

• Об’єднаний доступ до глобальної пам’яті: потоки варпа мають звертатися до послідовних адрес пам’яті (крок 1). Необ’єднаний доступ фрагментується на кілька транзакцій. За змоги забезпечуйте розкладку «структура масивів» (SoA) замість «масив структур» (AoS).
• Огрублення потоків: призначайте кілька елементів на потік, щоб зменшити накладні витрати на планування та поліпшити паралелізм на рівні інструкцій у межах потоку. Корисно, коли заповненість уже висока.
• Злиття ядер: об’єднуйте кілька обмежених пам’яттю ядер в одне, щоб уникнути повторних читань/записів глобальної пам’яті між ядрами. Може підвищити арифметичну інтенсивність до обчислювального даху.
• Тайлінг для множення матриць: завантажуйте тайли з глобальної пам’яті у спільну, обчислюйте часткові суми, накопичуйте в регістрах. Перетворює O(N³) глобальних байтів на O(N³/T) з розміром тайла T, різко підвищуючи арифметичну інтенсивність.
• Примітиви рівня варпа: використовуйте інструкції перемішування (__shfl_xor_sync) для обміну в межах варпа без спільної пам’яті — нульова затримка, без конфліктів банків.
• Асинхронні копіювання (Ampere+): memcpy_async перекриває завантаження даних з обчисленнями за допомогою програмного конвеєра (упереджене завантаження двох тайлів по черзі).
• Змішана точність: FP16 має 2× пропускну здатність, тензорні ядра BF16/TF32 дають 4–8× пропускної здатності для множення матриць. Використовуйте накопичення FP16/BF16 там, де точність дозволяє.