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

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

François Beaufort
François Beaufort

רקע

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

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

הדוגמה הבאה ממחישה איך לכתוב ארבעה בייטים של זיכרון מאגר נתונים זמני שאפשר לגשת אליו מה-GPU. הוא מבצע קריאה ל-device.createBuffer(), שלוקח את גודל המאגר והשימוש בו. למרות שסימון השימוש GPUBufferUsage.MAP_WRITE אינו נדרש עבור הקריאה הספציפית הזו, נבהיר שאנחנו רוצים לכתוב למאגר הזה. התוצאה היא אובייקט מאגר GPU שממופה בזמן היצירה, הודות לערך mappedAtCreation שמוגדר כ-True. לאחר מכן, אפשר לאחזר את מאגר הנתונים הבינארי של הנתונים הבינאריים המשויך של ה-GPU על ידי קריאה לשיטה getMappedRange() של מאגר הנתונים הבינארי של GPU.

כתיבת בייטים היא מוכרת אם כבר ניסיתם להשתמש ב-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.

תכנות הצללה

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

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

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

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

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

יצירת מאגרי GPU

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

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

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

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

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

בואו נראה איך תוכנת ההצללה (shader) הזאת תפעל ב-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().

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

בקוד שלנו, התוצאה שנרשמה בלוח 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.