תחילת העבודה עם GPU Compute באינטרנט

הפוסט הזה בוחן את ה-WebGPU API הניסיוני באמצעות דוגמאות ו להתחיל לבצע חישובים מקבילים של נתונים באמצעות ה-GPU.

François Beaufort
François Beaufort
GitHub

רקע

כפי שאולי כבר ידוע לך, יחידת העיבוד הגרפי (GPU) היא במחשב שהתמחה במקור בעיבוד גרפי. עם זאת, ב-10 השנים האחרונות, התפתחה שמאפשרת למפתחים להטמיע סוגים רבים של אלגוריתמים לעבד גרפיקה תלת-ממדית, תוך ניצול הארכיטקטורה הייחודית של GPU. היכולות האלה נקראות GPU Compute, והשימוש ב-GPU מעבד עזר למחשוב מדעי לשימוש כללי נקרא 'לשימוש כללי'. תכנות GPU (GPGPU).

GPU Compute תרם באופן משמעותי לפריחה האחרונה של למידת המכונה, רשתות נוירונים קונבולציה ומודלים אחרים יכולים לנצל כדי לפעול בצורה יעילה יותר במעבדי GPU. עם פלטפורמת האינטרנט הנוכחית ל-W3C אין יכולות מחשוב של GPU, ו-"GPU for the Web" קבוצת קהילה היא תכנון ממשק API שיחשוף את ממשקי ה-API המודרניים של GPU הזמינים ברוב מכשירים נוכחיים. ה-API הזה נקרא WebGPU.

WebGPU הוא ממשק API ברמה נמוכה, כמו WebGL. היא עוצמתית מאוד ודי מילונית, שתראו. אבל זה בסדר. מה שאנחנו מחפשים הוא ביצועים.

במאמר הזה אתמקד בחלק של GPU Compute ב-WebGPU, האמת היא שאנחנו רק מרקדים, כדי שתוכלו להתחיל לשחק שלו. אני אתעמק בנושא ואסביר על עיבוד של WebGPU (לוח הציור, טקסטורה, וכו') במאמרים הבאים.

גישה ל-GPU

הגישה ל-GPU קלה ב-WebGPU. אנחנו מתקשרים אל navigator.gpu.requestAdapter() מחזירה הבטחת JavaScript שתפתור באופן אסינכרוני עם GPU עם מתאם בלבד. המתאם הזה הוא כרטיס המסך. אפשר לשלב ביניהם (באותו צ'יפ כמו המעבד) או דיסקרטי (בדרך כלל כרטיס PCIe גדול יותר אך צורכת יותר חשמל).

אחרי שמוצאים את מתאם ה-GPU, צריך להתקשר אל adapter.requestDevice() כדי לקבל הבטחה שתפתרו את הבעיה במכשיר GPU שבו תשתמשו כדי לבצע חישובי GPU.

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

שתי הפונקציות משתמשות באפשרויות שמאפשרות לפרט את מתאם המתח (העדפת חשמל) ואת המכשיר (תוספים, מגבלות) הרצויים. עבור כדי לפשט את הפשטות, נשתמש באפשרויות ברירת המחדל במאמר הזה.

כתיבת זיכרון מאגר נתונים זמני

בואו נראה איך להשתמש ב-JavaScript כדי לכתוב נתונים בזיכרון עבור ה-GPU. הזה התהליך אינו פשוט בגלל המודל 'ארגז חול' שמשמש בדפדפנים אחרים.

הדוגמה הבאה מראה איך לכתוב ארבעה בייטים במאגר נתונים זמני ונגיש מה-GPU. הפונקציה קוראת לפונקציה device.createBuffer(), שגודלה מאגר הנתונים הזמני והשימוש בו. למרות שדגל השימוש GPUBufferUsage.MAP_WRITE הוא לא נדרשת לשיחה הספציפית הזו, בואו נבהיר שאנחנו רוצים לכתוב למאגר הנתונים הזמני. התוצאה היא אובייקט של מאגר נתונים זמני של GPU שממופה בזמן היצירה, mappedAtCreation מוגדר כ-True. לאחר מכן, מאגר הנתונים הבינאריים הגולמי המשויך יכול מאוחזר על ידי קריאה למאגר הנתונים הזמני של ה-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, נדרש דגל שימוש חדש GPUBufferUsage.COPY_SRC. השנייה מאגר הנתונים הזמני של ה-GPU נוצר במצב לא ממופה הפעם עם device.createBuffer() דגל השימוש שלו הוא 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, צריך לבצע קריאה ל-gpuReadBuffer.mapAsync() באמצעות GPUMapMode.READ. הוא מחזיר הבטחה שתפתור את הבעיה כשמאגר הנתונים הזמני של ה-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 שמוגדרת כ-True נקראת.

תכנות Shader

תוכנות שרצות ב-GPU שמבצעות רק חישובים (ולא שרטטו משולשים) נקראים תוכנות הצללה למחשוב. הן מבוצעות במקביל על ידי מאות של ליבות GPU (קטנות מליבות המעבד (CPU)) שפועלות יחד כדי לעבד . הקלט והפלט שלהם משמשים במאגרי נתונים זמניים ב-WebGPU.

כדי להמחיש את השימוש בתוכנות הצללה למחשוב ב-WebGPU, נשחק עם מטריצות אתיקה של כפל, אלגוריתם נפוץ בלמידת מכונה שמתואר בהמשך.

תרשים הכפל של המטריצה
תרשים הכפל של Matrix

בקצרה, זה מה שנעשה:

  1. יוצרים שלושה מאגרי נתונים זמניים של GPU (שניים כדי להכפיל את המטריצות ואחת בשביל מטריצת התוצאות)
  2. תיאור הקלט והפלט של תוכנת ההצללה למחשוב
  3. הידור של הקוד של תוכנת ההצללה (shader) המחשוב
  4. הגדרת צינור עיבוד נתונים של מחשוב
  5. שליחה באצווה של פקודות מקודדות ל-GPU
  6. קריאת מאגר הנתונים הזמני של מטריצת התוצאות של ה-GPU

יצירת מאגרי GPU

כדי לשמור על פשטות, המטריצות מיוצגות כרשימה של צפים ומספרי נקודות. הרכיב הראשון הוא מספר השורות, והרכיב השני מספר העמודות, והשאר הוא המספרים בפועל של המטריצה.

ייצוג פשוט של מטריצה ב-JavaScript ושל המקבילה שלה בסימון מתמטי
ייצוג פשוט של מטריצה ב-JavaScript והמקבילה שלה בסימון מתמטי

שלושת מאגרי הנתונים הזמניים של ה-GPU הם מאגרי אחסון זמניים, כי אנחנו צריכים לאחסן ולאחזר נתונים את תוכנת ההצללה למחשוב. זו הסיבה לכך שהסימונים לשימוש במאגר הנתונים הזמני של GPU כוללים GPUBufferUsage.STORAGE לכולם. דגל השימוש במטריצת התוצאות כולל גם GPUBufferUsage.COPY_SRC כי הוא יועתק למאגר נתונים אחר למשך קריאה לאחר שכל פקודות התור של 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. קישור פריסת קבוצה מגדירה את ממשק הקלט/פלט שנדרש על ידי תוכנת ההצללה, קבוצת קישור מייצגת את נתוני הקלט/פלט בפועל עבור תוכנת ההצללה.

בדוגמה הבאה, הפריסה של קבוצת הקישורים מצפה לשני מאגרי אחסון לקריאה בלבד קישורי רשומות ממוספרות 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)

קוד ההצללה למחשוב להכפלת מטריצות כתוב ב-WGSL, WebGPU Shader Language, שניתן לתרגום באופן טריוויאלי ל-SPIR-V. ללא בפירוט, אפשר לראות מתחת לשלושת מאגרי האחסון הזמניים שזוהו עם 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(). נדרשים שני ארגומנטים: פריסת קבוצת הקישור שיצרנו קודם, ו שלב המגדיר את נקודת הכניסה של תוכנת ההצללה (shader) שלנו (פונקציית WGSL main) והמודול של תוכנת ההצללה (shader) המחשוב בפועל שנוצר באמצעות device.createShaderModule().

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

שליחת פקודות

אחרי יצירת קבוצת קישור באמצעות שלושת מאגרי ה-GPU והמחשוב שלנו עם פריסת קבוצה של קשרים, זה הזמן להשתמש בהם.

בואו נתחיל מקודד של מעבר מחשוב שניתן לתכנות commandEncoder.beginComputePass() נשתמש בזה כדי לקודד פקודות GPU שתבצעו את ההכפלה של המטריצה. מגדירים את צינור עיבוד הנתונים באמצעות passEncoder.setPipeline(computePipeline) וקבוצת הקישור שלו באינדקס 0 עם passEncoder.setBindGroup(0, bindGroup). האינדקס 0 תואם עיטור group(0) בקוד WGSL.

עכשיו נדבר על האופן שבו תוכנת ההצללה למחשוב תרוץ על ה-GPU. שלנו הוא להפעיל את התוכנית הזו במקביל לכל תא במטריצת התוצאות, שלב אחרי שלב. למשל, כשמטריצת תוצאות בגודל 16 על 32 צריכה לקודד פקודת הביצוע, ב-@workgroup_size(8, 8), נקרא passEncoder.dispatchWorkgroups(2, 4) או passEncoder.dispatchWorkgroups(16 / 8, 32 / 8). הארגומנט הראשון 'x' הוא המאפיין הראשון, השני, 'y' הוא המאפיין השני, והפקודה האחרונה היא 'z' הוא המאפיין השלישי שברירת המחדל שלו היא 1, כי אנחנו לא צריכים אותו כאן. בעולם המחשוב של GPU, קידוד פקודה להפעלת פונקציית ליבה בקבוצת נתונים נקרא 'שליחת נתונים'.

ביצוע במקביל לכל תא מטריצת התוצאות
ביצוע במקביל לכל תא מטריצת התוצאות

גודל הרשת של קבוצת העבודה בתוכנת ההצללה (shader) המחשוב שלנו הוא (8, 8) ב-WGSL לכן, 'x' ו-y שהם בהתאמה מספר השורות של את המטריצה הראשונה ואת מספר העמודות במטריצה השנייה יחולקו עד 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 על ידי התקשרות device.queue.submit() מחליפים בפקודות ה-GPU.

// 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 ומחכים עד שההבטחה החוזרת תפתור את הבעיה, שמצביעה על כך מאגר הנתונים הזמני של ה-GPU ממופה עכשיו. בשלב הזה אפשר לקבל טווח עם gpuReadBuffer.getMappedRange().

תוצאת הכפלה של המטריצה
תוצאת הכפל של Matrix

בקוד שלנו, התוצאה שמתועדת במסוף JavaScript של כלי הפיתוח היא "2, 2, 50, 60, 114, 140 אינץ'.

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

מעולה! הצלחת. אפשר לשחק עם הדוגמה.

טריק אחד אחרון

אחת הדרכים להקל על קריאת הקוד היא להשתמש שיטת getBindGroupLayout בצינור עיבוד הנתונים של המחשוב כדי להסיק מהי קבוצת הקישור ממודול של תוכנת ההצללה. הטריק הזה מבטל את הצורך ליצור פריסת צינור עיבוד נתונים בהתאמה אישית וציון פריסת צינור עיבוד נתונים במחשוב כפי שאפשר לראות בהמשך.

איור של 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 על 256.

נקודת ההשוואה בין GPU למעבד (CPU)
בנצ'מרק למעבד (GPU) למעבד (CPU)

המאמר הזה היה רק תחילת המסע שלי עם התנסות ב-WebGPU. יש עוד ציפייה בקרוב נוסיף עוד ניתוחים מעמיקים לגבי GPU Compute והסבר על תהליך העיבוד. (בד ציור, טקסטורה, דגול) פועל ב-WebGPU.