เริ่มต้นใช้งาน GPU Compute บนเว็บ

โพสต์นี้จะสำรวจ WebGPU API เวอร์ชันทดลองผ่านตัวอย่างและช่วยให้คุณเริ่มต้นใช้งานการประมวลผลแบบขนานกันโดยใช้ GPU

François Beaufort
François Beaufort
GitHub

ข้อมูลเบื้องต้น

อย่างที่คุณอาจทราบแล้ว หน่วยประมวลผลกราฟิก (GPU) คือระบบย่อยอิเล็กทรอนิกส์ภายในคอมพิวเตอร์ที่เดิมทีออกแบบมาเพื่อประมวลผลกราฟิกโดยเฉพาะ อย่างไรก็ตาม ในช่วง 10 ปีที่ผ่านมา สถาปัตยกรรมของ GPU ได้พัฒนาไปให้มีความยืดหยุ่นมากขึ้น ซึ่งช่วยให้นักพัฒนาซอฟต์แวร์สามารถใช้อัลกอริทึมได้หลายประเภท ไม่ใช่แค่แสดงผลกราฟิก 3 มิติเท่านั้น และยังใช้ประโยชน์จากสถาปัตยกรรมเฉพาะของ GPU ได้อีกด้วย ความสามารถเหล่านี้เรียกว่า GPU Compute และการใช้ GPU เป็นหน่วยประมวลผลร่วมสำหรับการประมวลผลทางวิทยาศาสตร์ทั่วไปเรียกว่าการเขียนโปรแกรม GPU (GPGPU) ทั่วไป

การคำนวณด้วย GPU มีส่วนสำคัญอย่างยิ่งในการเติบโตของแมชชีนเลิร์นนิงในช่วงที่ผ่านมา เนื่องจากโครงข่ายประสาทแบบ Convolution และโมเดลอื่นๆ ใช้ประโยชน์จากสถาปัตยกรรมนี้เพื่อทำงานได้อย่างมีประสิทธิภาพมากขึ้นใน GPU เนื่องจากแพลตฟอร์มเว็บในปัจจุบันไม่มีความสามารถในการประมวลผลด้วย GPU กลุ่มชุมชน "GPU สําหรับเว็บ" ของ W3C จึงออกแบบ API เพื่อแสดง GPU API สมัยใหม่ที่มีให้บริการในอุปกรณ์ส่วนใหญ่ในปัจจุบัน API นี้เรียกว่า WebGPU

WebGPU เป็น API ระดับล่าง เช่น WebGL ซึ่งมีประสิทธิภาพมากและค่อนข้างจะแสดงผลข้อมูลอย่างละเอียดตามที่คุณจะได้เห็น แต่ไม่เป็นไร สิ่งที่เรามองหาคือประสิทธิภาพ

ในบทความนี้ เราจะมุ่งเน้นที่ส่วน GPU Compute ของ WebGPU และขอบอกตามตรงว่าเราเพิ่งจะเริ่มต้นเท่านั้น เพื่อให้คุณเริ่มเล่นด้วยตนเองได้ เราจะเจาะลึกและอธิบายการเรนเดอร์ WebGPU (Canvas, พื้นผิว และอื่นๆ) ในบทความที่กำลังจะเผยแพร่

เข้าถึง GPU

การเข้าถึง GPU ใน WebGPU ทำได้ง่าย การเรียก navigator.gpu.requestAdapter() จะแสดงผลพรอมต์ JavaScript ที่จะแก้ไขแบบไม่พร้อมกันด้วยอะแดปเตอร์ GPU ให้คิดว่าอะแดปเตอร์นี้เป็นการ์ดกราฟิก โดยอาจเป็นแบบรวม (อยู่ในชิปเดียวกับ CPU) หรือแบบแยก (โดยปกติจะเป็นการ์ด PCIe ที่มีประสิทธิภาพมากกว่าแต่ใช้พลังงานมากกว่า)

เมื่อคุณมีแอปแดปเตอร์ GPU แล้ว ให้เรียกใช้ adapter.requestDevice() เพื่อรับสัญญาว่าจะดำเนินการซึ่งจะแสดงผลเป็นอุปกรณ์ GPU ที่คุณจะใช้ทำการคำนวณบางอย่างของ GPU

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();

ฟังก์ชันทั้ง 2 ฟังก์ชันใช้ตัวเลือกที่ให้คุณระบุประเภทของอะแดปเตอร์ (ค่ากำหนดกำลังไฟ) และอุปกรณ์ (ส่วนขยาย ข้อจำกัด) ที่ต้องการได้ เราจะใช้ตัวเลือกเริ่มต้นในบทความนี้เพื่อให้เข้าใจง่าย

เขียนหน่วยความจำบัฟเฟอร์

มาดูวิธีใช้ JavaScript เพื่อเขียนข้อมูลลงในหน่วยความจําสําหรับ GPU กระบวนการนี้ไม่ตรงไปตรงมาเนื่องจากรูปแบบแซนด์บ็อกซ์ที่ใช้ในเว็บเบราว์เซอร์สมัยใหม่

ตัวอย่างด้านล่างแสดงวิธีเขียน 4 ไบต์ไปยังบัฟเฟอร์หน่วยความจำที่เข้าถึงได้จาก GPU ซึ่งจะเรียก device.createBuffer() ซึ่งจะรับขนาดบัฟเฟอร์และการใช้งาน แม้ว่าการเรียกใช้นี้ไม่จำเป็นต้องใช้ Flag การใช้งาน GPUBufferUsage.MAP_WRITE แต่เราต้องการเขียนลงในบัฟเฟอร์นี้อย่างชัดเจน การดำเนินการนี้ส่งผลให้ออบเจ็กต์บัฟเฟอร์ GPU แมปเมื่อสร้างเนื่องจากมีการตั้งค่า mappedAtCreation เป็น "จริง" จากนั้นเรียกใช้เมธอดบัฟเฟอร์ GPU getMappedRange() เพื่อเรียกข้อมูลไบนารีดิบที่เกี่ยวข้อง

การเขียนไบต์เป็นสิ่งที่คุณคุ้นเคยหากเคยเล่นกับ ArrayBuffer มาก่อน ให้ใช้ TypedArray แล้วคัดลอกค่าลงไป

// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuBuffer = device.createBuffer({
  mappedAtCreation: true,
  size: 4,
  usage: GPUBufferUsage.MAP_WRITE
});
const arrayBuffer = gpuBuffer.getMappedRange();

// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);

เมื่อถึงจุดนี้ ระบบจะแมปบัฟเฟอร์ GPU ซึ่งหมายความว่าบัฟเฟอร์เป็นของ CPU และเข้าถึงได้แบบอ่าน/เขียนจาก JavaScript คุณต้องยกเลิกการแมปเพื่อให้ GPU เข้าถึงได้ ซึ่งทำได้ง่ายๆ เพียงเรียกใช้ gpuBuffer.unmap()

แนวคิดของการแมป/ไม่ได้แมปจําเป็นเพื่อป้องกันเงื่อนไขการแข่งขันที่ GPU และ CPU เข้าถึงหน่วยความจําพร้อมกัน

อ่านหน่วยความจำบัฟเฟอร์

มาดูวิธีคัดลอกบัฟเฟอร์ GPU ไปยังบัฟเฟอร์ GPU อื่นและอ่านกลับกัน

เนื่องจากเราเขียนลงในบัฟเฟอร์ GPU แรกและต้องการคัดลอกไปยังบัฟเฟอร์ GPU ที่ 2 จึงต้องใช้ Flag การใช้งานใหม่ GPUBufferUsage.COPY_SRC ระบบจะสร้างบัฟเฟอร์ GPU ตัวที่ 2 ในสถานะไม่ได้แมปในครั้งนี้ด้วย device.createBuffer() Flag การใช้งานคือ GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ เนื่องจากจะใช้เป็นปลายทางของบัฟเฟอร์ GPU แรกและอ่านใน JavaScript เมื่อมีการเรียกใช้คำสั่งการคัดลอก GPU

// Get a GPU buffer in a mapped state and an arrayBuffer for writing.
const gpuWriteBuffer = device.createBuffer({
  mappedAtCreation: true,
  size: 4,
  usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.COPY_SRC
});
const arrayBuffer = gpuWriteBuffer.getMappedRange();

// Write bytes to buffer.
new Uint8Array(arrayBuffer).set([0, 1, 2, 3]);

// Unmap buffer so that it can be used later for copy.
gpuWriteBuffer.unmap();

// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
  size: 4,
  usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});

เนื่องจาก GPU เป็นหน่วยประมวลผลร่วมอิสระ ระบบจึงจะเรียกใช้คําสั่ง GPU ทั้งหมดแบบไม่พร้อมกัน ด้วยเหตุนี้จึงมีรายการคําสั่ง GPU ที่สร้างขึ้นและส่งเป็นกลุ่มเมื่อจําเป็น ใน WebGPU ตัวแปลงรหัสคำสั่ง GPU ที่ device.createCommandEncoder() แสดงผลคือออบเจ็กต์ JavaScript ที่สร้างชุดคำสั่ง "บัฟเฟอร์" ที่จะส่งไปยัง GPU ณ เวลาหนึ่งๆ ในทางกลับกัน เมธอดใน GPUBuffer จะเป็น "ไม่มีการบัฟเฟอร์" ซึ่งหมายความว่าจะดำเนินการแบบอะตอม ณ เวลาที่มีการเรียกใช้

เมื่อคุณมีโปรแกรมเข้ารหัสคําสั่ง GPU แล้ว ให้เรียกใช้ copyEncoder.copyBufferToBuffer() ตามที่แสดงด้านล่างเพื่อเพิ่มคําสั่งนี้ลงในคิวคําสั่งสําหรับการดําเนินการในภายหลัง สุดท้าย ให้เสร็จสิ้นการเข้ารหัสคําสั่งโดยเรียกใช้ copyEncoder.finish() และส่งคําสั่งเหล่านั้นไปยังคิวคําสั่งของอุปกรณ์ GPU คิวมีหน้าที่จัดการการส่งผ่าน device.queue.submit() ที่มีคำสั่ง GPU เป็นอาร์กิวเมนต์ ซึ่งจะดำเนินการคำสั่งทั้งหมดที่จัดเก็บไว้ในอาร์เรย์ตามลำดับ

// Encode commands for copying buffer to buffer.
const copyEncoder = device.createCommandEncoder();
copyEncoder.copyBufferToBuffer(
  gpuWriteBuffer /* source buffer */,
  0 /* source offset */,
  gpuReadBuffer /* destination buffer */,
  0 /* destination offset */,
  4 /* size */
);

// Submit copy commands.
const copyCommands = copyEncoder.finish();
device.queue.submit([copyCommands]);

ณ จุดนี้ ระบบได้ส่งคําสั่งคิว GPU แล้ว แต่ไม่จําเป็นต้องดําเนินการ หากต้องการอ่านบัฟเฟอร์ GPU ที่ 2 ให้เรียกใช้ gpuReadBuffer.mapAsync() ด้วย GPUMapMode.READ โดยจะแสดงผลลัพธ์เป็น Promise ที่จะแสดงผลเมื่อมีการแมปบัฟเฟอร์ GPU จากนั้นรับช่วงที่แมปด้วย gpuReadBuffer.getMappedRange() ซึ่งมีค่าเดียวกับบัฟเฟอร์ GPU แรกเมื่อมีการเรียกใช้คําสั่ง GPU ที่รอดำเนินการทั้งหมดแล้ว

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const copyArrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Uint8Array(copyArrayBuffer));

คุณสามารถลองใช้ตัวอย่างนี้

กล่าวโดยย่อ สิ่งที่คุณต้องจำเกี่ยวกับการดำเนินการกับหน่วยความจำบัฟเฟอร์มีดังนี้

  • คุณต้องยกเลิกการแมปบัฟเฟอร์ GPU เพื่อใช้ในการส่งคิวอุปกรณ์
  • เมื่อแมปแล้ว คุณจะอ่านและเขียนบัฟเฟอร์ GPU ใน JavaScript ได้
  • ระบบจะแมปบัฟเฟอร์ GPU เมื่อมีการเรียกใช้ mapAsync() และ createBuffer() โดยตั้งค่า mappedAtCreation เป็น "จริง"

โปรแกรมการเขียนเฉดสี

โปรแกรมที่ทำงานบน GPU ซึ่งทำการประมวลผลเท่านั้น (และไม่วาดสามเหลี่ยม) เรียกว่า คอมพิวตเชดเดอร์ โดยระบบจะประมวลผลไปพร้อมกันโดยใช้แกน GPU หลายร้อยแกน (ซึ่งเล็กกว่าแกน CPU) ที่ทำงานร่วมกันเพื่อประมวลผลข้อมูล อินพุตและเอาต์พุตของ WebGPU คือบัฟเฟอร์

เราจะอธิบายการใช้เชดเดอร์การประมวลผลใน WebGPU โดยใช้การคูณเมทริกซ์ ซึ่งเป็นอัลกอริทึมทั่วไปในแมชชีนเลิร์นนิงดังที่แสดงด้านล่าง

แผนภาพการคูณเมทริกซ์
แผนภาพการคูณเมทริกซ์

โดยสรุปแล้ว สิ่งที่เราจะทำมีดังนี้

  1. สร้างบัฟเฟอร์ GPU 3 รายการ (2 รายการสำหรับเมทริกซ์ที่จะคูณ และ 1 รายการสำหรับเมทริกซ์ผลลัพธ์)
  2. อธิบายอินพุตและเอาต์พุตสำหรับเชดเดอร์การประมวลผล
  3. คอมไพล์โค้ด Shader การคำนวณ
  4. ตั้งค่าไปป์ไลน์การประมวลผล
  5. ส่งคำสั่งที่เข้ารหัสไปยัง GPU แบบเป็นกลุ่ม
  6. อ่านบัฟเฟอร์ GPU ของเมทริกซ์ผลลัพธ์

การสร้างบัฟเฟอร์ GPU

เราจะแสดงเมทริกซ์เป็นรายการตัวเลขทศนิยมเพื่อความสะดวก องค์ประกอบแรกคือจํานวนแถว องค์ประกอบที่ 2 คือจํานวนคอลัมน์ และที่เหลือคือจํานวนจริงของเมทริกซ์

การนําเสนอเมทริกซ์อย่างง่ายใน JavaScript และเทียบเท่าในเครื่องหมายทางคณิตศาสตร์
การนําเสนอเมทริกซ์อย่างง่ายใน JavaScript และเทียบเท่าในสัญลักษณ์ทางคณิตศาสตร์

บัฟเฟอร์ GPU 3 รายการคือบัฟเฟอร์พื้นที่เก็บข้อมูลเนื่องจากเราต้องจัดเก็บและเรียกข้อมูลในคอมพิวตเชดเดอร์ ด้วยเหตุนี้ แฟล็กการใช้งานบัฟเฟอร์ GPU จึงมี GPUBufferUsage.STORAGE สำหรับรายการทั้งหมด นอกจากนี้ Flag การใช้งานเมทริกซ์ผลลัพธ์ยังมี GPUBufferUsage.COPY_SRC ด้วย เนื่องจากระบบจะคัดลอก Flag นี้ไปยังบัฟเฟอร์อื่นเพื่ออ่านเมื่อระบบเรียกใช้คิว GPU ทั้งหมดแล้ว

const adapter = await navigator.gpu.requestAdapter();
if (!adapter) { return; }
const device = await adapter.requestDevice();


// First Matrix

const firstMatrix = new Float32Array([
  2 /* rows */, 4 /* columns */,
  1, 2, 3, 4,
  5, 6, 7, 8
]);

const gpuBufferFirstMatrix = device.createBuffer({
  mappedAtCreation: true,
  size: firstMatrix.byteLength,
  usage: GPUBufferUsage.STORAGE,
});
const arrayBufferFirstMatrix = gpuBufferFirstMatrix.getMappedRange();
new Float32Array(arrayBufferFirstMatrix).set(firstMatrix);
gpuBufferFirstMatrix.unmap();


// Second Matrix

const secondMatrix = new Float32Array([
  4 /* rows */, 2 /* columns */,
  1, 2,
  3, 4,
  5, 6,
  7, 8
]);

const gpuBufferSecondMatrix = device.createBuffer({
  mappedAtCreation: true,
  size: secondMatrix.byteLength,
  usage: GPUBufferUsage.STORAGE,
});
const arrayBufferSecondMatrix = gpuBufferSecondMatrix.getMappedRange();
new Float32Array(arrayBufferSecondMatrix).set(secondMatrix);
gpuBufferSecondMatrix.unmap();


// Result Matrix

const resultMatrixBufferSize = Float32Array.BYTES_PER_ELEMENT * (2 + firstMatrix[0] * secondMatrix[1]);
const resultMatrixBuffer = device.createBuffer({
  size: resultMatrixBufferSize,
  usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC
});

เชื่อมโยงเลย์เอาต์กลุ่มและเชื่อมโยงกลุ่ม

แนวคิดของเลย์เอาต์กลุ่มการเชื่อมโยงและกลุ่มการเชื่อมโยงมีไว้สำหรับ WebGPU โดยเฉพาะ เลย์เอาต์กลุ่มการเชื่อมโยงจะกำหนดอินเทอร์เฟซอินพุต/เอาต์พุตที่ชิดเดอร์คาดหวัง ส่วนกลุ่มการเชื่อมโยงจะแสดงข้อมูลอินพุต/เอาต์พุตจริงของชิดเดอร์

ในตัวอย่างนี้ เลย์เอาต์กลุ่มการเชื่อมโยงจะคาดหวังว่าจะมีบัฟเฟอร์พื้นที่เก็บข้อมูลแบบอ่านอย่างเดียว 2 รายการที่การเชื่อมโยงรายการที่ระบุหมายเลข 0, 1 และบัฟเฟอร์พื้นที่เก็บข้อมูลที่ 2 สำหรับเชดเดอร์การประมวลผล กลุ่มการเชื่อมโยงที่กำหนดไว้สำหรับเลย์เอาต์กลุ่มการเชื่อมโยงนี้จะเชื่อมโยงบัฟเฟอร์ GPU กับรายการต่างๆ ดังนี้ gpuBufferFirstMatrix กับรายการการเชื่อมโยง 0, gpuBufferSecondMatrix กับรายการการเชื่อมโยง 1 และ resultMatrixBuffer กับรายการการเชื่อมโยง 2

const bindGroupLayout = device.createBindGroupLayout({
  entries: [
    {
      binding: 0,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {
        type: "read-only-storage"
      }
    },
    {
      binding: 1,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {
        type: "read-only-storage"
      }
    },
    {
      binding: 2,
      visibility: GPUShaderStage.COMPUTE,
      buffer: {
        type: "storage"
      }
    }
  ]
});

const bindGroup = device.createBindGroup({
  layout: bindGroupLayout,
  entries: [
    {
      binding: 0,
      resource: {
        buffer: gpuBufferFirstMatrix
      }
    },
    {
      binding: 1,
      resource: {
        buffer: gpuBufferSecondMatrix
      }
    },
    {
      binding: 2,
      resource: {
        buffer: resultMatrixBuffer
      }
    }
  ]
});

โค้ด Shader ของ Compute

โค้ด Shader แบบประมวลผลสำหรับการคูณเมทริกซ์เขียนด้วย WGSL ซึ่งเป็นภาษา Shader ของ WebGPU ที่แปลเป็น SPIR-V ได้โดยง่าย โดยไม่ต้องลงรายละเอียด คุณควรเห็นบัฟเฟอร์พื้นที่เก็บข้อมูล 3 รายการที่ระบุด้วย var<storage> ด้านล่าง โปรแกรมจะใช้ firstMatrix และ secondMatrix เป็นอินพุต และ resultMatrix เป็นเอาต์พุต

โปรดทราบว่าบัฟเฟอร์พื้นที่เก็บข้อมูลแต่ละรายการมีการตกแต่ง binding ที่ใช้ซึ่งสอดคล้องกับดัชนีเดียวกันที่กําหนดไว้ในเลย์เอาต์กลุ่มการเชื่อมโยงและกลุ่มการเชื่อมโยงที่ประกาศไว้ด้านบน

const shaderModule = device.createShaderModule({
  code: `
    struct Matrix {
      size : vec2f,
      numbers: array<f32>,
    }

    @group(0) @binding(0) var<storage, read> firstMatrix : Matrix;
    @group(0) @binding(1) var<storage, read> secondMatrix : Matrix;
    @group(0) @binding(2) var<storage, read_write> resultMatrix : Matrix;

    @compute @workgroup_size(8, 8)
    fn main(@builtin(global_invocation_id) global_id : vec3u) {
      // Guard against out-of-bounds work group sizes
      if (global_id.x >= u32(firstMatrix.size.x) || global_id.y >= u32(secondMatrix.size.y)) {
        return;
      }

      resultMatrix.size = vec2(firstMatrix.size.x, secondMatrix.size.y);

      let resultCell = vec2(global_id.x, global_id.y);
      var result = 0.0;
      for (var i = 0u; i < u32(firstMatrix.size.y); i = i + 1u) {
        let a = i + resultCell.x * u32(firstMatrix.size.y);
        let b = resultCell.y + i * u32(secondMatrix.size.y);
        result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b];
      }

      let index = resultCell.y + resultCell.x * u32(secondMatrix.size.y);
      resultMatrix.numbers[index] = result;
    }
  `
});

การตั้งค่าไปป์ไลน์

พอร์ทัลการคำนวณคือออบเจ็กต์ที่อธิบายการดำเนินการประมวลผลที่เราจะดำเนินการ สร้างโดยโทรไปที่ device.createComputePipeline() โดยจะใช้อาร์กิวเมนต์ 2 รายการ ได้แก่ เลย์เอาต์กลุ่มการเชื่อมโยงที่เราสร้างขึ้นก่อนหน้านี้ และระยะการประมวลผลที่กําหนดจุดแรกเข้าของ Shader การประมวลผล (ฟังก์ชัน main WGSL) และโมดูล Shader การประมวลผลจริงที่สร้างด้วย device.createShaderModule()

const computePipeline = device.createComputePipeline({
  layout: device.createPipelineLayout({
    bindGroupLayouts: [bindGroupLayout]
  }),
  compute: {
    module: shaderModule,
    entryPoint: "main"
  }
});

การส่งคําสั่ง

หลังจากสร้างอินสแตนซ์ของกลุ่มการเชื่อมโยงด้วยบัฟเฟอร์ GPU 3 รายการและไปป์ไลน์การประมวลผลที่มีเลย์เอาต์กลุ่มการเชื่อมโยงแล้ว ก็ถึงเวลาใช้งาน

มาเริ่มโปรแกรมเปลี่ยนไฟล์พาสการประมวลผลแบบกำหนดโปรแกรมได้ด้วย commandEncoder.beginComputePass() เราจะใช้สิ่งนี้เพื่อเข้ารหัสคําสั่ง GPU ที่จะทำการดำเนินการคูณเมทริกซ์ ตั้งค่าไปป์ไลน์ด้วย passEncoder.setPipeline(computePipeline) และกลุ่มการเชื่อมโยงที่ดัชนี 0 ด้วย passEncoder.setBindGroup(0, bindGroup) ดัชนี 0 สอดคล้องกับการตกแต่ง group(0) ในโค้ด WGSL

ทีนี้มาพูดถึงวิธีที่คอมพิวตเชดเดอร์นี้จะทำงานบน GPU เป้าหมายของเราคือเรียกใช้โปรแกรมนี้แบบขนานกันสำหรับแต่ละเซลล์ของเมทริกซ์ผลลัพธ์ทีละขั้นตอน ตัวอย่างเช่น สำหรับเมทริกซ์ผลลัพธ์ขนาด 16 x 32 หากต้องการเข้ารหัสคำสั่งการดําเนินการใน @workgroup_size(8, 8) เราจะเรียก passEncoder.dispatchWorkgroups(2, 4) หรือ passEncoder.dispatchWorkgroups(16 / 8, 32 / 8) โดยอาร์กิวเมนต์แรก "x" คือมิติข้อมูลแรก อาร์กิวเมนต์ที่ 2 "y" คือมิติข้อมูลที่สอง และอาร์กิวเมนต์สุดท้าย "z" คือมิติข้อมูลที่สามซึ่งมีค่าเริ่มต้นเป็น 1 เนื่องจากเราไม่ได้ใช้มิติข้อมูลนี้ ในการคำนวณของ GPU การเข้ารหัสคำสั่งเพื่อเรียกใช้ฟังก์ชันเคอร์เนลในชุดข้อมูลเรียกว่าการจัดเตรียม

การดำเนินการพร้อมกันสำหรับเซลล์เมทริกซ์ผลลัพธ์แต่ละเซลล์
การดําเนินการพร้อมกันสําหรับเซลล์เมทริกซ์ผลลัพธ์แต่ละเซลล์

ขนาดตารางกริดของกลุ่มงานสำหรับนักคำนวณของเราคือ (8, 8) ในโค้ด WGSL ด้วยเหตุนี้ "x" และ "y" ซึ่งเป็นจำนวนแถวของเมทริกซ์แรกและจำนวนคอลัมน์ของเมทริกซ์ที่ 2 ตามลำดับจึงต้องหารด้วย 8 ตอนนี้เราส่งการเรียกใช้การประมวลผลด้วย passEncoder.dispatchWorkgroups(firstMatrix[0] / 8, secondMatrix[1] / 8) ได้แล้ว จํานวนตารางกริดของกลุ่มงานที่จะเรียกใช้คืออาร์กิวเมนต์ dispatchWorkgroups()

ดังที่เห็นในภาพด้านบน แต่ละชิดเดอร์จะมีสิทธิ์เข้าถึงbuiltin(global_invocation_id)ออบเจ็กต์ที่ไม่ซ้ำกันซึ่งจะใช้เพื่อดูว่าต้องคำนวณเซลล์ผลลัพธ์ใดในเมทริกซ์

const commandEncoder = device.createCommandEncoder();

const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, bindGroup);
const workgroupCountX = Math.ceil(firstMatrix[0] / 8);
const workgroupCountY = Math.ceil(secondMatrix[1] / 8);
passEncoder.dispatchWorkgroups(workgroupCountX, workgroupCountY);
passEncoder.end();

หากต้องการสิ้นสุดโปรแกรมเปลี่ยนไฟล์พาสคอมพิวต์ ให้โทรหา passEncoder.end() จากนั้นสร้างบัฟเฟอร์ GPU เพื่อใช้เป็นปลายทางในการคัดลอกบัฟเฟอร์เมทริกซ์ผลลัพธ์ด้วย copyBufferToBuffer สุดท้าย ให้เข้ารหัสคําสั่งด้วย copyEncoder.finish() และส่งคําสั่งเหล่านั้นไปยังคิวอุปกรณ์ GPU โดยเรียกใช้ copyEncoder.finish() ด้วยคําสั่ง GPUdevice.queue.submit()

// Get a GPU buffer for reading in an unmapped state.
const gpuReadBuffer = device.createBuffer({
  size: resultMatrixBufferSize,
  usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
});

// Encode commands for copying buffer to buffer.
commandEncoder.copyBufferToBuffer(
  resultMatrixBuffer /* source buffer */,
  0 /* source offset */,
  gpuReadBuffer /* destination buffer */,
  0 /* destination offset */,
  resultMatrixBufferSize /* size */
);

// Submit GPU commands.
const gpuCommands = commandEncoder.finish();
device.queue.submit([gpuCommands]);

อ่านเมทริกซ์ผลลัพธ์

การอ่านเมทริกซ์ผลลัพธ์นั้นง่ายเพียงเรียกใช้ gpuReadBuffer.mapAsync() ด้วย GPUMapMode.READ แล้วรอให้ Promise ที่แสดงผลแสดงผล ซึ่งบ่งบอกว่าตอนนี้มีการแมปบัฟเฟอร์ GPU แล้ว เมื่อถึงจุดนี้ คุณสามารถรับช่วงที่มีการแมปด้วย gpuReadBuffer.getMappedRange()

ผลคูณเมทริกซ์
ผลลัพธ์ของการคูณเมทริกซ์

ในโค้ดของเรา ผลลัพธ์ที่บันทึกไว้ในคอนโซล JavaScript ของ DevTools คือ "2, 2, 50, 60, 114, 140"

// Read buffer.
await gpuReadBuffer.mapAsync(GPUMapMode.READ);
const arrayBuffer = gpuReadBuffer.getMappedRange();
console.log(new Float32Array(arrayBuffer));

ยินดีด้วย คุณทำสำเร็จแล้ว คุณสามารถเล่นกับตัวอย่างเพลงได้

เคล็ดลับสุดท้าย

วิธีหนึ่งในการทําให้โค้ดอ่านง่ายขึ้นคือการใช้getBindGroupLayoutเมธอดที่มีประโยชน์ของไปป์ไลน์การประมวลผลเพื่ออนุมานเลย์เอาต์กลุ่มการเชื่อมโยงจากโมดูล Shader เคล็ดลับนี้ช่วยให้คุณไม่ต้องสร้างเลย์เอาต์กลุ่มการเชื่อมโยงที่กำหนดเองและระบุเลย์เอาต์ไปป์ไลน์ในไปป์ไลน์การประมวลผลตามที่แสดงด้านล่าง

มีภาพ getBindGroupLayout สำหรับตัวอย่างก่อนหน้า

 const computePipeline = device.createComputePipeline({
-  layout: device.createPipelineLayout({
-    bindGroupLayouts: [bindGroupLayout]
-  }),
   compute: {

-// Bind group layout and bind group
- const bindGroupLayout = device.createBindGroupLayout({
-   entries: [
-     {
-       binding: 0,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "read-only-storage"
-       }
-     },
-     {
-       binding: 1,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "read-only-storage"
-       }
-     },
-     {
-       binding: 2,
-       visibility: GPUShaderStage.COMPUTE,
-       buffer: {
-         type: "storage"
-       }
-     }
-   ]
- });
+// Bind group
  const bindGroup = device.createBindGroup({
-  layout: bindGroupLayout,
+  layout: computePipeline.getBindGroupLayout(0 /* index */),
   entries: [

ผลการค้นหาเกี่ยวกับประสิทธิภาพ

ดังนั้น การทำงานคูณเมทริกซ์ใน GPU เปรียบเทียบกับการทำงานใน CPU อย่างไร เราจึงเขียนโปรแกรมที่อธิบายไปสำหรับ CPU เพื่อหาคำตอบ และดังที่คุณเห็นในกราฟด้านล่าง การใช้พลังของ GPU อย่างเต็มรูปแบบดูเหมือนจะเป็นตัวเลือกที่ชัดเจนเมื่อขนาดของเมทริกซ์มากกว่า 256 x 256

การเปรียบเทียบ GPU กับ CPU
การเปรียบเทียบประสิทธิภาพ GPU กับ CPU

บทความนี้เป็นเพียงจุดเริ่มต้นของเส้นทางการสำรวจ WebGPU โปรดรอบทความเพิ่มเติมในเร็วๆ นี้ซึ่งจะเจาะลึกเรื่อง GPU Compute และวิธีการทำงานของการแสดงผล (Canvas, Texture, Sampler) ใน WebGPU