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

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

François Beaufort
François Beaufort
GitHub

רקע

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

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

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

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

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

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

עכשיו נראה איך מעתיקים מאגר נתונים של 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 ומבצעות רק חישובים (ולא מציירות משולשים) נקראות compute shaders. הן מתבצעות במקביל על ידי מאות ליבות GPU (שהן קטנות יותר מליבות CPU) שפועלות יחד כדי לנתח את הנתונים. הקלט והפלט שלהם הם מאגרים ב-WebGPU.

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

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

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

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

יצירת מאגרים של GPU

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

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

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

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

קוד של שדה חישוב (compute shader)

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

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

שליחת פקודות

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

נתחיל להפעיל מקודד של כרטיס Compute Pass שניתן לתכנות באמצעות 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, קידוד של פקודה להפעלת פונקציית ליבה על קבוצת נתונים נקרא 'הקצאה'.

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

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

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

בקוד שלנו, התוצאה שמתועדת במסוף 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 ועל האופן שבו עיבוד (קנבס, טקסטורה, Sampler) פועל ב-WebGPU.