Learn the fundamentals of General-Purpose GPU (GPGPU) compute pipelines, grid threading models (Workgroups, local and global invocation IDs), resource synchronization via barriers, and dynamic Storage Buffer allocation in WebGPU.

This lesson is currently only available in Vietnamese. Please switch the language toggle in the menu to Vietnamese to read the full guide and take the interactive quiz.

Trong các bài học trước, chúng ta đã tập trung hoàn toàn vào Render Pipeline – tức là quy trình đưa dữ liệu đỉnh qua các bộ lọc ánh sáng, bóng đổ rồi vẽ chúng trực tiếp ra màn hình canvas. Tuy nhiên, sức mạnh thực sự của GPU hiện đại nằm ở khả năng tính toán song song tổng quát (GPGPU - General-Purpose computing on GPUs).

WebGPU hỗ trợ tính toán song song hiệu năng cao thông qua Compute Pipeline. Trong bài học này, chúng ta sẽ đi sâu vào mô hình phân luồng Compute Shader, cách trao đổi dữ liệu qua Storage Buffer và thực hiện đo lường hiệu năng (Benchmark) của phép nhân ma trận vuông khổng lồ giữa CPU (JavaScript) và GPU (WebGPU Compute).

1. Khái niệm GPGPU & Compute Pipeline

Khác với CPU chỉ sở hữu từ 4 đến 64 lõi tính toán phức tạp chạy tuần tự tối ưu, GPU sở hữu hàng ngàn lõi tính toán đơn giản chạy song song theo kiến trúc SIMT (Single Instruction, Multiple Threads). Quy trình xử lý Compute Pipeline của WebGPU cực kỳ tối giản vì nó không cần qua giai đoạn lắp ráp đỉnh (vertex assembly) hay rasterization của màn hình vẽ. Nhiệm vụ duy nhất của nó là nhận dữ liệu thô, tính toán song song, và trả lại kết quả thô.

2. Mô hình phân luồng: Workgroups & Invocation IDs

Khi dispatch một Compute Shader, chúng ta không định nghĩa các luồng chạy riêng rẽ mà chia chúng thành các nhóm công việc gọi là Workgroup (Nhóm làm việc). Mỗi Workgroup lại chứa một lưới đa chiều các luồng tính toán (threads).

Để định vị từng luồng tính toán đơn lẻ trong một mạng lưới khổng lồ, WGSL cung cấp các định danh tích hợp sẵn (System Built-ins):

• @builtin(local_invocation_id): Tọa độ 3D của luồng nằm trong Workgroup hiện tại (chạy từ (0,0,0) đến (X-1, Y-1, Z-1)).
• @builtin(global_invocation_id): Tọa độ 3D tuyệt đối của luồng trên toàn bộ lưới tính toán tổng thể. Đây là ID quan trọng nhất để ánh xạ luồng vào phần tử mảng dữ liệu.
• @builtin(workgroup_id): Tọa độ định danh của chính nhóm Workgroup đó trên lưới Grid.

3. Bộ nhớ chia sẻ (Shared Memory) & Barrier

Bên cạnh Storage Buffer (nằm ở VRAM toàn cục của GPU, truy cập chậm hơn), các luồng nằm trong cùng một Workgroup có thể chia sẻ dữ liệu tốc độ cực nhanh thông qua bộ đệm cục bộ Shared Memory (khai báo bằng từ khóa var<workgroup> trong WGSL).

Vì các luồng chạy song song bất đồng bộ, để tránh xung đột đọc/ghi bộ nhớ chia sẻ này (Data Race), ta phải sử dụng lệnh đồng bộ workgroupBarrier(). Lệnh này hoạt động giống như một rào cản: tất cả các luồng trong nhóm bắt buộc phải dừng lại đợi cho đến khi toàn bộ luồng khác cùng nhóm chạy tới dòng này, đảm bảo dữ liệu ghi vào Shared Memory đã hoàn tất trước khi luồng kế tiếp đọc ra.

4. Thuật toán nhân ma trận song song trên GPU

Phép nhân hai ma trận vuông $A \times B = C$ kích thước $N \times N$ thông thường trên CPU chạy 3 vòng lặp lồng nhau mất độ phức tạp thời gian cực lớn là \(O(N^3)\).

Với GPU Compute Shader, ta có thể song song hóa hoàn toàn: Ta dispatch một lưới gồm $N \times N$ luồng song song. Mỗi luồng tương ứng với tọa độ hàng row (trục Y) và cột col (trục X) của ma trận kết quả $C$. Luồng này chỉ chạy duy nhất 1 vòng lặp chiều dài $N$ để tính tổng tích chập hàng-cột của phần tử $C[row, col]$ đó, giảm độ phức tạp thời gian chạy song song lý thuyết xuống chỉ còn \(O(N)\)!

Sân chơi tương tác: CPU vs GPU Matrix Benchmark

Dưới đây là một phòng thí nghiệm hiệu năng thực tế. Hãy chọn kích thước ma trận vuông $N \times N$, tạo ngẫu nhiên dữ liệu ma trận, và chạy so sánh tốc độ thực thi giữa CPU (vòng lặp JavaScript thô) và GPU (WebGPU Compute Shader):

struct Dims {
  size: u32,
};

@group(0) @binding(0) var<uniform> u: Dims;
@group(0) @binding(1) var<storage, read> a: array<f32>;
@group(0) @binding(2) var<storage, read> b: array<f32>;
@group(0) @binding(3) var<storage, read_write> c: array<f32>;

@compute @workgroup_size(16, 16)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
  let size = u.size;
  let row = id.y;
  let col = id.x;

  // Tránh ghi đè nếu lưới Workgroup vượt quá kích thước ma trận thực tế
  if (row >= size || col >= size) {
    return;
  }

  var sum: f32 = 0.0;
  for (var k: u32 = 0u; k < size; k = k + 1u) {
    sum = sum + a[row * size + k] * b[k * size + col];
  }

  c[row * size + col] = sum;
}

// 1. Khai báo Compute Pipeline
const pipeline = device.createComputePipeline({
  layout: 'auto',
  compute: { module, entryPoint: 'main' }
});

// 2. Dispatch phân luồng song song trên GPU
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(pipeline);
passEncoder.setBindGroup(0, bindGroup);
// Chia nhỏ công việc thành các Workgroups kích cỡ 16x16
const workgroupsX = Math.ceil(size / 16);
const workgroupsY = Math.ceil(size / 16);
passEncoder.dispatchWorkgroups(workgroupsX, workgroupsY, 1);
passEncoder.end();

// 3. Sao chép kết quả về CPU qua Staging Buffer
commandEncoder.copyBufferToBuffer(gpuBufC, 0, gpuStaging, 0, byteSize);
device.queue.submit([commandEncoder.finish()]);

// Đợi và map dữ liệu về bộ nhớ CPU đọc được
await gpuStaging.mapAsync(GPUMapMode.READ);
const result = new Float32Array(gpuStaging.getMappedRange().slice());
gpuStaging.unmap();