שיפורים ב-WebAssembly וב-WebGPU לשמירה מהירה יותר על ה-Web AI, חלק {/9}2

המסמך הזה הוא המשך של השיפורים ב-WebAssembly וב-WebGPU לפיצ'רים מהירים יותר של Web AI, חלק 1. אנחנו ממליצים לקרוא את הפוסט הזה או לצפות בשיחה ב-IO 24 לפני שממשיכים.

Austin Eng
Austin Eng
Deepti Gandluri
Deepti Gandluri
François Beaufort
François Beaufort

WebGPU

באמצעות WebGPU אפליקציות אינטרנט מאפשרות לאפליקציות אינטרנט לגשת לחומרת ה-GPU של הלקוח כדי לבצע חישוב יעיל ומקביל מאוד. מאז ההשקה של WebGPU ב-Chrome, קיבלנו הדגמות מדהימות של בינה מלאכותית (AI) ולמידת מכונה (ML) באינטרנט.

לדוגמה, טכנולוגיית הדיפוזיה היציבה באינטרנט הראו שאפשר להשתמש ב-AI כדי ליצור תמונות מטקסט ישירות בדפדפן. מוקדם יותר השנה, צוות Mediapipe של Google עצמו פרסם תמיכה ניסיונית בהסקת מסקנות ממודל שפה גדול (LLM).

האנימציה הבאה מציגה את Gemma, מודל השפה הגדול (LLM) של Google בקוד פתוח, שפועל רק במכשיר ב-Chrome בזמן אמת.

ההדגמה הבאה של הפנים החיבוק של Meta Anything Model מייצרת מסכות באיכות גבוהה לאובייקטים.

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

בנצ'מרק של WebGPU להטמעת טקסט של התכונה 'חיבוק פנים', המצולם במהירויות גבוהות מאוד בהשוואה להטמעה של מעבד (CPU) באותו מודל. במחשב נייד של Apple M1 Max, WebGPU מהיר פי 30. אחרים דיווחו כי WebGPU מאיץ את נקודת ההשוואה יותר מ-120 פעמים.

שיפור תכונות WebGPU ל-AI ולמידת מכונה

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

בין השיפורים הרבים ב-WebGPU בשנה האחרונה, המשכנו להוסיף עוד יכולות לשיפור הביצועים של למידת מכונה ו-AI באינטרנט. לאחרונה השקנו שתי תכונות חדשות: נקודה צפה (floating-point) של 16 ביט ומוצרים ארוזים בנקודות מספרים שלמים.

נקודה צפה (floating-point) של 16 ביט

חשוב לזכור שעומסי עבודה של למידת מכונה לא דורשים דיוק. shader-f16 היא תכונה שמאפשרת שימוש בסוג f16 בשפת ההצללה של WebGPU. סוג הנקודה הצפה הזה תופס 16 ביט, במקום את 32 הביטים הרגילים. ל-f16 יש טווח קטן יותר והוא פחות מדויק, אבל להרבה מודלים של למידת מכונה זה מספיק.

התכונה הזו מגבירה את היעילות בכמה דרכים:

  • צמצום הזיכרון: טנזרים עם רכיבי f16 תופסים חצי משטח האחסון, ולכן השימוש בזיכרון יתקצר בחצי. לעיתים קרובות, חישובי GPU עוברים צוואר בקבוק ברוחב הפס של הזיכרון, כך שחצי מהזיכרון יכול בדרך כלל לגרום לתוכנות הצללה לפעול מהר פי שניים. מבחינה טכנית אין צורך ב-f16 כדי לחסוך ברוחב הפס של הזיכרון. ניתן לאחסן את הנתונים בפורמט ברמת דיוק נמוכה ולאחר מכן להרחיב אותם ל-f32 באופן מלא בכלי ההצללה לצורך חישוב. עם זאת, ה-GPU משקיע יותר כוח מחשוב כדי לארוז ולפרק את הנתונים.

  • צמצום של המרות הנתונים: f16 משתמש בפחות מחשוב על ידי צמצום המרות הנתונים. ניתן לאחסן נתונים ברמת דיוק נמוכה ולהשתמש בהם ישירות ללא המרה.

  • קיבולת מקבילה מוגדלת: מעבדי GPU מודרניים יכולים להכניס יותר ערכים בו-זמנית ביחידות הביצוע של ה-GPU, וכך לבצע מספר גדול יותר של חישובים מקבילים. לדוגמה, יחידת GPU שתומכת בעד 5 טריליון פעולות נקודה צפה (floating-point) של f32 לשנייה עשויה לתמוך ב-10 טריליון פעולות נקודה צפה (floating-point) לשנייה מסוג f16.

צילום מסך של בנצ'מרק ל-WebGPU להטמעת טקסט
עם shader-f16, נקודת ההשוואה WebGPU להטמעת טקסט של פועלת בנצ'מרק פי 3 יותר מהר מ-f32 במחשב נייד Apple M1 Max.

WebLLM הוא פרויקט שיכול להריץ כמה מודלים גדולים של שפה. היא משתמשת ב-Apache TVM, מסגרת מהדר (compiler) בקוד פתוח ללמידת מכונה.

ביקשתי מ-WebLLM לתכנן נסיעה לפריז באמצעות המודל לאמה 3 של שמונה מיליארד פרמטרים. התוצאות מראות שבמהלך שלב המילוי מראש של המודל, f16 מהיר פי 2.1 מאשר f32. בשלב הפענוח הוא מהיר יותר פי 1.3.

קודם צריך לוודא שהמתאם ל-GPU תומך ב-f16. אם הוא זמין, צריך להפעיל אותו באופן מפורש בבקשת מכשיר GPU. אם f16 לא נתמך, לא ניתן לבקש אותו במערך requiredFeatures.

// main.js

const adapter = await navigator.gpu.requestAdapter();
const supportsF16 = adapter.features.has('shader-f16');
if (supportsF16) {
  // Use f16.
  const device = await adapter.requestDevice({
    requiredFeatures: ['shader-f16'],
  });
  initApp(device);
}

לאחר מכן, בכלי ההצללה של WebGPU, צריך להפעיל באופן מפורש את f16 בחלק העליון. לאחר מכן אפשר להשתמש בו בכלי ההצללה כמו בכל סוג אחר של נתונים צפים.

// my-shader.wgsl

enable f16;

struct Data {
  values : array<vec4<f16>>
}
@group(0) @binding(0) var<storage, read> data : Data;
@compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid : vec3u) {
  let value : vec4<f16> = data.values[gid.x];
  ...
}

מוצרי נקודות שלמים ארוזים

הרבה מודלים עדיין פועלים היטב עם 8 ביטים של דיוק (חצי מ-f16). שיטה זו פופולרית בקרב מודלי שפה גדולים (LLM) ובמודלים של תמונות לצורך פילוח וזיהוי אובייקטים. עם זאת, איכות הפלט של המודלים פוחתת ברמת דיוק נמוכה יותר, לכן כימות ב-8 ביט לא מתאים לכל אפליקציה.

מעט מאוד מעבדי GPU תומכים במקור של ערכי 8 ביט. כאן נכנסים למוצרים ארוזים של נקודות עם מספרים שלמים. שלחנו את DP4a ב-Chrome 123.

למעבדי GPU מודרניים יש הוראות מיוחדות לשליפת שני מספרים שלמים של 32 ביט, לפירוש כל אחד מהם כ-4 מספרים שלמים בחבילה של 8 ביט ולחישוב המכפלה בין הרכיבים שלהם.

האפשרות הזו שימושית במיוחד ל-AI וללמידת מכונה, כי ליבות הכפל של מטריצות מורכבות מהרבה מאוד תוצרי נקודות.

לדוגמה, נכפיל מטריצה של 4 x 8 בווקטור של 8 x 1. כדי לחשב את הערך הזה, צריך לקחת 4 נקודות כדי לחשב כל אחד מהערכים בווקטור הפלט. א', ב', ג' ו-ד'.

דוגמה להכפלה של מטריצה וקטורית

תהליך החישוב של כל אחד מהפלטים האלה זהה; נבחן את השלבים הנדרשים כדי לחשב אחד מהם. לפני כל חישוב, עלינו להמיר תחילה את נתוני המספרים השלמים ב-8 ביט לסוג שבו ניתן לבצע חשבון, כמו f16. לאחר מכן נריץ הכפלה ברמת הרכיבים, ולבסוף נחבר את כל המכפלות. בסך הכול, בכל ההכפלה של המטריצה והוקטורים, אנחנו מבצעים 40 המרות מסוג מספרים שלמים כדי לפרוק את הנתונים, 32 הכפלות צפות ו-28 הוספות צפות.

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

עבור כל אחד מהפלטים בווקטור התוצאה, אנחנו מבצעים שתי פעולות מוצר באמצעות נקודה ארוזה באמצעות dot4U8Packed המובנה של WebGPU Shading Language, ואז מחברים את התוצאות. לסה"כ, בכל ההכפלה של המטריצה בווקטור, לא מבצעים המרת נתונים. אנחנו מפעילים 8 מוצרי נקודות ארוזים ו-4 הוספות של מספרים שלמים.

דוגמה להכפלה של מטריצה-וקטור של מספרים שלמים ארוזים

בדקנו מוצרים ארוזים של מספרים שלמים עם נתונים של 8 ביט במגוון מעבדי GPU לצרכנים. בהשוואה לנקודה צפה (floating-point) של 16 ביט, אפשר לראות שקצב 8 ביט מהיר פי 1.6 עד 2.8 ביט. כשאנחנו משתמשים גם במוצרי נקודות שלמים ארוזים, הביצועים טובים עוד יותר. הוא מהיר פי 1.7 עד פי 2.9.

צילום מסך של האצת הכפלת וקטורים של מטריצות: f16 לעומת u8
תרשים 1: האצת הווקטור של מטריצות, השוואה בין f16 ל-U8 ו-U8 עם נקודה4U8Packed.

אפשר לבדוק אם הדפדפן תומך בעזרת המאפיין wgslLanguageFeatures. אם ה-GPU לא תומך במקור במוצרי נקודות ארוזים, הדפדפן משלים את ההטמעה של הדפדפן שלו.

// main.js

if (navigator.gpu.wgslLanguageFeatures.has('packed_4x8_integer_dot_product')) {
  // Use dot4U8Packed, dot4I8Packed builtin
  // functions in the shaders.
}

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

לפני – תוכנת ההצללה של WebGPU שמצטברת מוצרי נקודות חלקיים במשתנה 'sum'. בסוף הלולאה, 'sum' מחזיק את מכפלת הנקודות המלאה בין וקטור לבין שורה אחת של מטריצת הקלט.

// my-dot-product.wgsl

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
  var sum : f16;
  let start = gid.x * uniforms.dim;
  for (var i = 0u; i < uniforms.dim; i++) {
    let v1 : vec4<f16> = vector.values[i];
    let v2 : vec4<f16> = matrix.values[start + i];
    sum += dot(v1, v2);
  }
}

אחרי — כלי ההצללה של WebGPU שנכתב כדי להשתמש במוצרי נקודות שלמים ארוזים. ההבדל העיקרי הוא שבמקום לטעון 4 ערכים צפים מחוץ לווקטור ומהמטריצה, תוכנת ההצללה הזו טוענת מספר שלם יחיד של 32 ביט. המספר השלם הזה בפורמט 32 ביט מכיל את הנתונים של ארבעה ערכים של מספרים שלמים של 8 ביט. אחר כך קוראים לפונקציה dot4U8Packed כדי לחשב את המכפלה של שני הערכים.

// my-dot-product.wgsl

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid : vec3u) {
  var sum : f32;
  let start = gid.x * uniforms.dim;
  for (var i = 0u; i < uniforms.dim; i++) {
    let v1 : u32 = vector.values[i];
    let v2 : u32 = matrix.values[start + i];
    sum += dot4U8Packed(v1, v2);
  }
}

התכונות של Chrome עם נקודה צפה (floating-point) בחבילה עם מספרים שלמים ארוזות הן התכונות שמאיצות את יכולות ה-AI ולמידת המכונה (ML). נקודה צפה (floating-point) של 16 ביט זמינה כשהחומרה תומכת בה, ו-Chrome מטמיע בכל המכשירים מוצרים ארוזים בנקודות מספרים שלמים.

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

תכונות מוצעות

בהמשך, אנחנו בודקים שתי תכונות נוספות: תת-קבוצות ומטריצה שיתופית.

התכונה של תת-קבוצות מאפשרת מקבילות ברמת ה-SIMD לתקשר או לבצע פעולות מתמטיות קולקטיביות, כמו סכום ליותר מ-16 מספרים. כך ניתן לשתף נתונים ביעילות בין שרשורים. תת-קבוצות נתמכות בממשקי API מודרניים של מעבדי GPU, עם שמות שונים ובצורות מעט שונות.

זקקנו את הקבוצה המשותפת להצעה שהעברנו לקבוצת התקינה של WebGPU. בנוסף, יצרנו אב טיפוס של קבוצות משנה ב-Chrome מאחורי דגל ניסיוני, והוספנו את התוצאות הראשוניות שלנו לדיון. הבעיה העיקרית היא איך להבטיח התנהגות ניידת.

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

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

כדי להעריך את הביצועים של פעולות תת-קבוצות, ביישום אמיתי, שילבנו תמיכה ניסיונית בקבוצות משנה ב-MediaPipe ובדקנו אותה עם אב הטיפוס של Chrome לפעולות משנה.

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

צילום מסך של האצת המהירות של תתי-קבוצות בהסקת מסקנות LLM מ-MediaPipe
תרשים 2. קבוצות המשנה מזרזות פי 2.5 את הרצת המילוי מראש ב-GPU של Intel Tiger Lake GT2, עם תמיכה ניסיונית ב-Chrome וב-Mediapipe.

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

צילום מסך של האצת תת-קבוצה להכפלת מטריצה
תרשים 3. החלת תת-קבוצות על הכפלת מטריצות עשויה לשפר את הביצועים עוד יותר.

קשה לבצע אופטימיזציה של ה-GPU

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

אתם רוצים לצמצם את רוחב הפס של הזיכרון תוך שימוש מלא בשרשורי המחשוב של ה-GPU.

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

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

אפשר לראות שלזיכרון יש השפעה משמעותית, לפעמים היא משפיעה אפילו יותר מאשר תת-קבוצות. ב-GPU 6, swizzle מספק מהירות גבוהה פי 12 ואילו תתי-קבוצות מספקות מהירות פי 13. במשולב, מהירות הטעינה שלהם היא אדירה פי 26. במקרים של מעבדי GPU אחרים, לפעמים הביצועים של החלקה ותתי-קבוצות מניבים ביצועים טובים יותר מאשר כל אחת מהם בנפרד. ובמעבדי GPU אחרים, רק שימוש ב-swizzle מניב את הביצועים הטובים ביותר.

צילום מסך של האצת אסטרטגיות להכפלת מטריצות
תרשים 4.

לפעמים צריך הרבה ניסיון כדי לכוונן ולבצע אופטימיזציה של האלגוריתמים של ה-GPU כדי שיפעלו היטב על כל חלק בחומרה. אבל למרבה המזל, יש עבודה מושקעת עצומה של עבודה מושקעת במסגרות של ספריות ברמה גבוהה יותר, כמו Mediapipe, Transformers.js, Apache TVM, ONNX Runtime Web ועוד.

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

חטיפות דסקית

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

המטרה שלנו היא למקסם את היכולות של הפלטפורמה תוך שמירה על היתרונות של האינטרנט: פוטנציאל החשיפה, השימושיות והניידות. אנחנו לא עושים את זה לבד. אנחנו עובדים בשיתוף פעולה עם ספקי דפדפנים אחרים ב-W3C ושותפי פיתוח רבים.

אנחנו מקווים שתזכרו את הדברים הבאים, כשאתם עובדים עם WebAssembly ו-WebGPU:

  • אפשרות ההסקה מ-AI זמינה עכשיו באינטרנט בכל המכשירים. כך אפשר ליהנות מהיתרון של ההרצה במכשירים של הלקוחות, כמו עלות השרת מופחתת, זמן אחזור קצר והגברת הפרטיות.
  • חלק גדול מהתכונות שעליהן מדברים רלוונטיות בעיקר למחברי ה-framework, אבל האפליקציות שלכם יכולות להפיק תועלת ללא תקורה משמעותית.
  • הסטנדרטים באינטרנט הם גמישים ומתפתחים, ואנחנו תמיד מחפשים משוב. משתפים את שלכם ב-WebAssembly וב-WebGPU.

אישורים

רצינו להודות לצוות הגרפיקה באינטרנט של Intel, שהשתתף באופן משמעותי בפיתוח WebGPU f16 ובתכונות המוצר של הנקודות השלמות. אנחנו רוצים להודות לשאר החברים בקבוצות העבודה WebAssembly ו-WebGPU ב-W3C, כולל ספקי הדפדפנים האחרים.

תודה לצוותי ה-AI והלמידת מכונה ב-Google ובקהילת הקוד הפתוח, שהם שותפים מדהימים. וכמובן, כל חברי הצוות שלנו מאפשרים את כל זה.