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

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

François Beaufort

רקע

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

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

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

עכשיו נראה איך מעתיקים מאגר 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() עם הערך true של mappedAtCreation.

תכנות של Shader

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

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

בקצרה, אלה הפעולות שנבצע:

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

יצירת מאגרי GPU

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

שלושת המאגרים של ה-GPU הם מאגרי אחסון, כי אנחנו צריכים לאחסן ולאחזר נתונים ב-compute shader. זו הסיבה לכך שדגלי השימוש במאגר של ה-GPU כוללים את הערך GPUBufferUsage.STORAGE עבור כולם. גם לדגל השימוש במטריצת התוצאות יש GPUBufferUsage.COPY_SRC כי הוא יועתק למאגר אחר לצורך קריאה אחרי שכל הפקודות בתור של המעבד הגרפי יבוצעו.

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 לחישוב

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

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

המאמר הזה הוא רק ההתחלה של המסע שלי בנושא WebGPU. בקרוב יפורסמו מאמרים נוספים עם ניתוחים מעמיקים יותר של GPU Compute ושל אופן הפעולה של rendering (קנבס, טקסטורה, sampler) ב-WebGPU.