לדלג לתוכן

1.5 Tensor Cores והאצת מטריצות הרצאה

בשיעור 1.3 פירקנו את ה-SM (מעבד רב-זרמי - Streaming Multiprocessor) לרכיביו וראינו את ה-CUDA Core (ליבת CUDA) - יחידת ה-ALU הסקלרית שמבצעת הוראה אחת על thread (thread) אחד, מתוזמנת ב-warp-ים (קבוצות של 32 threads) על ידי ארבעת ה-warp schedulers של ה-SM. אבל אם תסתכלו על גיליון הנתונים של H100 ותראו את המספר המסחרר של כמעט אלף TFLOPS, תגלו שרובו המכריע כלל לא מגיע מה-CUDA Cores. הוא מגיע מיחידת חישוב שנייה, שונה במהותה, שיושבת לצד ה-CUDA Cores בכל SM - ה-Tensor Core (ליבת טנזור). בשיעור הזה נכיר את המנוע שמניע את מהפכת הבינה המלאכותית: ליבה שמבצעת בהוראה אחת כפל של מטריצות שלמות, לא של סקלרים. נבין למה היא כמו מעבד CISC בעולם של RISC, נפרק הוראת HMMA אחת עד לרמת ה-MAC הבודד, נראה איך warp שלם משתף פעולה כדי לחשב tile אחד, ונכיר את חומרת התמיכה שנבנתה סביבה - warpgroup, ה-TMA וזיכרון הטנזור. את איך לתכנת את כל זה בפועל, דרך ספריות כמו cuBLAS ו-CUTLASS, נלמד בפרק 6; כאן אנחנו מבינים את החומרה.

מהי ליבת טנזור - Tensor Core

ליבת טנזור היא, בהגדרה של ה-glossary, ליבת GPU ש"פועלת על מטריצות שלמות בכל הוראה". זו הנקודה שצריך לתפוס לפני הכל: בעוד ש-CUDA Core מקבלת שני סקלרים ומחזירה סקלר, ליבת טנזור מקבלת tiles של מטריצות שלמים ומבצעת עליהם פעולה אחת גדולה - כפל-וצבירה של מטריצות (matrix multiply-accumulate, בקיצור MMA / MAC) מהצורה:

D = A * B + C

כאשר A, B, C, D הם מטריצות (ליתר דיוק, tiles של מטריצות). ההוראה מכפילה את A ב-B, מוסיפה את מטריצת הצבירה C, וכותבת ל-D. זו בדיוק הפעולה שבלב כל שכבה של רשת נוירונים ובלב כל GEMM (General Matrix Multiply), ולכן זו הפעולה שאותה NVIDIA בחרה להאיץ בחומרה ייעודית.

RISC מול CISC - האנלוגיה המרכזית. ה-glossary ממסגר את ליבת הטנזור כחומרה בסגנון CISC, בניגוד ל-CUDA Core שהיא בסגנון RISC. ליבת CUDA היא כמו הוראת RISC: הוראה פשוטה, אחידה, שעושה מעט עבודה (חיבור, כפל, fma). ליבת טנזור היא כמו הוראת CISC "שמנה": הוראה אחת שצורכת ומייצרת tiles של מטריצות שלמים ומבצעת אלפי פעולות כפל-וחיבור. מכאן נובע גם יתרון ההספק שלה, שנפרט מיד: היא פורשת את עלות שליפת ההוראה על פני הרבה מאוד אריתמטיקה.

המספרים הפיזיים ב-H100. בניגוד לגודש ה-CUDA Cores (128 יחידות FP32 ל-SM, כפי שראינו בשיעור 1.3), ליבות הטנזור הן גדולות פיזית ומעטות בהרבה:

מאפיין CUDA Core Tensor Core
סגנון RISC - הוראה פשוטה על סקלר CISC - הוראה על מטריצה שלמה
פעולה סקלר: d = a*b + c מטריצה: D = A*B + C
כמות ב-SM (H100) 128 יחידות FP32 4 בלבד - אחת לכל warp scheduler
גודל פיזי קטנה, רבות גדולה, מעטות
תפוקה יחסית (FLOPS) בסיס כלל אצבע: פי ~100

שימו לב לסימטריה היפה: ב-H100 יש ארבע ליבות טנזור ל-SM - בדיוק אחת לכל warp scheduler (זכרו מ-1.3 שלכל SM ארבע מחיצות תזמון, כל אחת עם ה-warp scheduler שלה). כל warp scheduler מזין את ליבת הטנזור "שלו". ההספק העצום - כלל האצבע של ה-glossary הוא פי כ-100 מ-CUDA Core - מושג לא בעזרת יותר יחידות, אלא בעזרת יחידות שכל אחת מהן עושה הרבה יותר עבודה בכל הוראה.

ציר זמן. ליבות הטנזור הוצגו לראשונה בדור Volta (הכרטיס V100), ומאז התפתחו בכל דור - Turing, Ampere (A100), Hopper (H100) ו-Blackwell (B200). בכל דור הן קיבלו רמות דיוק חדשות (FP16, BF16, TF32, INT8, FP8, ובבלקוול אף FP4) ותמיכה בחומרה סביבתית שנפגוש בהמשך השיעור.

אנטומיה של הוראה - HMMA16.16816.F32

בואו נפרק הוראת ליבת טנזור אחת, אמיתית, עד לרמת החישוב הבודד. ההוראה שנעבוד עליה היא HMMA16.16816.F32, ופענוח השם מלמד הכל:

HMMA 16 . 16 8 16 . F32
 |    |     |  | |    |
 |    |     |  | |    +-- accumulation precision: FP32 (single precision)
 |    |     m  n k        tile dimensions: m=16, n=8, k=16
 |    +-- input precision: FP16 (16-bit half precision)
 +-- Half-precision Matrix Multiply-Accumulate

כלומר: הקלטים A ו-B הם ב-דיוק חצי (FP16, 16 ביט), הצבירה מתבצעת ב-דיוק יחיד (FP32, 32 ביט float), וממדי הtile הם m=16, n=8, k=16. הפרדת דיוק הקלט מדיוק הצבירה היא קריטית: מכפילים במספרים "זולים" בני 16 ביט, אבל צוברים במצטבר "יקר" ומדויק בן 32 ביט, כדי שהשגיאה המספרית לא תתפוצץ לאורך k המכפלות.

כמה חישוב יש בהוראה אחת? מכפלת מטריצות A(m x k) * B(k x n) דורשת m * n * k פעולות כפל-וצבירה (MAC). נציב:

MACs להוראה = m * n * k = 16 * 8 * 16 = 2,048 MAC

כל MAC הוא כפל אחד ועוד חיבור אחד, כלומר שתי פעולות נקודה-צפה (FLOPs):

FLOPs להוראה = 2 * 2,048 = 4,096 FLOP

וכמה נופל על כל thread? זכרו שההוראה מבוצעת על ידי warp שלם - 32 threads. מחלקים:

MACs ל-thread = 2,048 / 32 = 64 MAC לכל thread

עצרו על המספר הזה. הוראה אחת שכל thread מנפיק גורמת ל-64 פעולות כפל-וצבירה לכל thread. השוו זאת ל-CUDA Core, שבה הוראת FFMA אחת = MAC אחד לכל thread. יחס של 64:1 בכמות האריתמטיקה לכל הוראה - זהו מקור היעילות.

מדוע זה חסכוני באנרגיה - הטיעון של CISC. ה-glossary מנסח זאת כך: היעילות באה מ"פעולה על יותר נתונים עבור שליפת הוראה אחת". כל הוראה עולה: יש לשלוף אותה מהcache, לפענח אותה, לתזמן אותה. בעולם RISC של ה-CUDA Cores, העלות הקבועה הזו נפרשת על MAC בודד. בעולם CISC של ליבת הטנזור, אותה עלות בדיוק נפרשת על 2,048 MAC. הפחתת הoverhead היחסית של שליפה-ופענוח היא בדיוק מה שמאפשר את פי ~100 בתפוקה ובחסכון האנרגטי.

ביצוע שיתופי ברמת ה-warp - warp-level execution

כאן שוכן המלכוד הרעיוני הגדול ביותר של ליבות הטנזור, וחובה להבין אותו נכון. thread בודד שמנפיק הוראת HMMA אחת אינו מחשב מטריצה שלמה בעצמו. במקום זאת, warp שלם משתף פעולה בביצוע הוראת המטריצה: הtile כולו פרוש על פני קובצי האוגרים (register files) של 32 הthreads גם יחד. אף thread לא מחזיק את המטריצה כולה; כל thread מחזיק פרוסה קטנה ממנה באוגרים שלו.

זו בדיוק הסיבה שבגללה ה-API של ליבות הטנזור מכונה wmma - Warp-level Matrix Multiply-Accumulate. ה"warp-level" בשם אינו קישוט; הוא אומר שיחידת התכנות הבסיסית כאן היא ה-warp, לא ה-thread. כל 32 הthreads חייבים להגיע להוראת ה-wmma יחד, מכונסים (converged), ולתרום את פרוסתם.

       naive illusion                         reality: cooperative warp

   (wrong - do not think this way)        (32 threads share one tile)

   thread 0 --> whole matrix D?          +-----------------------------+
                                          |  tile D, size 16 x 8        |
   no! no single thread                   |                             |
   holds the whole matrix.                 |  spread across the registers|
                                          |  of the warp's 32 threads    |
                                          |  (each thread: 64 MAC)        |
                                          +-----------------------------+
                                            t0  t1  t2 ... t30 t31
                                            each contributes a small slice

ההשלכה המעשית: אי אפשר "לפזר" עבודת ליבת טנזור על threads בודדים כמו שמפזרים חיבור וקטורים. אם warp מתפצל (warp divergence - התפצלות ה-warp, כפי שראינו בשיעור על ה-warp), הוראת ה-wmma פשוט לא תוכל לרוץ נכון. ליבת הטנזור היא מטבעה מבנה שיתופי ברמת ה-warp, ובדורות החדשים אף ברמת ה-warpgroup (ארבעה warps, 128 threads) - כפי שנראה בהמשך.

שלוש רמות התכנות - WMMA, PTX, SASS

זוכרים מ-0.4 ומ-1.3 שקוד GPU חי בשלוש רמות - CUDA C++ למעלה, PTX (Parallel Thread eXecution) באמצע כשפת ביניים וירטואלית, ו-SASS (Streaming ASSembler) למטה כאסמבלי מקורי. הוראת ליבת טנזור עוברת את אותו המסלול בדיוק, וכדאי לראות אותו מקצה לקצה.

רמה 1 - CUDA C++ (ה-WMMA API). בקוד גבוה כותבים אינטרינזיק אחד:

#include <mma.h>
using namespace nvcuda;
// ... inside the kernel, after loading fragments a, b and zeroing c ...
wmma::mma_sync(c, a, b, c);   // C = A*B + C at the warp level

רמה 2 - PTX (אינטרינזיק MMA וירטואלי). ה-mma_sync הבודד מתקמפל להוראת PTX וירטואלית אחת ברמת ה-warp:

wmma.mma.sync.aligned.col.row.m16n16k16.f32.f32

הפענוח: wmma.mma.sync = כפל-מטריצות מסונכרן ברמת warp; .col.row = פריסת הזיכרון של האופרנדים (column-major ו-row-major); .m16n16k16 = ממדי הtile הלוגי; .f32.f32 = דיוק הצבירה והפלט.

רמה 3 - SASS (הוראות HMMA מרובות). וכאן ההפתעה החשובה: הוראת wmma אחת ברמת PTX אינה מתקמפלת להוראת SASS אחת. היא מתפרקת לכמה הוראות HMMA קטנות יותר:

HMMA.1688.F32 R20, R12, R11, RZ
HMMA.1688.F32 R24, R12, R17, RZ
HMMA.1688.F32 R20, R14, R16, R20
HMMA.1688.F32 R24, R14, R18, R24

כל הוראת HMMA.1688.F32 היא tile m16 n8 k8, כלומר 16 * 8 * 8 = 1,024 MAC. ארבע מהן מכסות יחד 4 * 1,024 = 4,096 MAC - וזה בדיוק גודל ה-wmma הלוגי m16n16k16 (כי 16 * 16 * 16 = 4,096). כך החומרה מרכיבה tile לוגי גדול מכמה tiles קטנים של חומרה. הסיומת .F32 היא דיוק הצבירה, ו-RZ הוא אוגר האפס (Zero Register) - כאן הוא משמש כאשר מתחילים לצבור מאפס.

המיקרו-ארכיטקטורה סגורה. איך ליבת הטנזור בנויה בפנים? זה קניין רשום ולא מתועד. ההשערה הרווחת היא שמדובר במערך סיסטולי (systolic array) - רשת של יחידות כפל-וחיבור שדוחפות נתונים זו לזו בקצב שעון - אבל ה-glossary מדגיש במפורש שאין "קונסנזוס בספרות ה-microbenchmarking" על התכן הפנימי המדויק. מבחינתנו כמתכנתים, ליבת הטנזור היא קופסה שחורה שמקבלת tiles ומחזירה tiles; מה שקורה בפנים נשאר אצל NVIDIA.

אל תכתבו ליבות טנזור ביד - cuBLAS, CUTLASS, CuTe

אחרי כל הפירוק הזה, ההמלצה המעשית של ה-glossary חד-משמעית: אל תכתבו הוראות ליבת טנזור ביד. התכנות הישיר מול ה-WMMA API, ובוודאי מול PTX בכתב יד, הוא קשה, שביר, ורגיש מאוד לפריסת הזיכרון, ל-bank conflicts (התנגשויות בנקים) ב-shared memory, ולתזמון. קשה מאוד להגיע ביד לתפוקת השיא של החומרה. יתרה מזו, ב-Hopper וב-Blackwell השגת הביצועים המקסימליים מחייבת אינטרינזיקים של PTX - קוד CUDA C++ פשוט כבר אינו מספיק כדי לרווות את ליבות הטנזור.

לכן משתמשים בספריות שכתבו המומחים של NVIDIA:

ספרייה רמה מתי משתמשים
cuBLAS גבוהה, "קופסה שחורה" כפל מטריצות סטנדרטי (GEMM), קריאה אחת
CUTLASS תבניות C++ ניתנות להרכבה GEMM מותאם, היתוך (fusion) עם פעולות אחרות
CuTe DSL שפה ייעודית לתיאור tiles ופריסות שליטה עדינה בתנועת נתונים ובtiles

הערת מלכוד שתלווה אתכם לאורך כל פרק 6: cuBLAS היא column-major (סדר עמודות, בירושה מ-Fortran ו-BLAS הקלאסי), בעוד ש-C/C++ הם row-major (סדר שורות). זהו מקור באגים קלאסי - מטריצה שנראית נכונה ב-C תתפרש הפוך ב-cuBLAS. את הפתרון המעשי לזה (התעלול B*A במקום A*B, או שימוש ב-CUBLAS_OP_T) נראה בפרק 6. כאן חשוב רק שתזכרו: אתם קוראים לליבות הטנזור דרך ספרייה, לא כותבים להן הוראות. השיעור הזה נותן לכם את המודל המנטלי כדי להבין מה הספרייה עושה מתחת למכסה המנוע - ולמה, כשתפרופלו GEMM ותראו HMMA ב-SASS, תדעו בדיוק מה אתם מסתכלים עליו.

דיוקים נתמכים - supported precisions

הבחירה בין דיוק הקלט לדיוק הצבירה, שראינו בהוראת ה-HMMA (קלט FP16, צבירה FP32), אינה מקרית - היא הלב של תכנון ליבות הטנזור. הרעיון המרכזי: מכפילים בדיוק נמוך וזול, אבל צוברים בדיוק גבוה ומדויק. הכפל בדיוק נמוך חוסך שטח סיליקון ואנרגיה ומכפיל את התפוקה, בעוד הצבירה בדיוק גבוה מונעת מהשגיאה המספרית להצטבר לאורך k המכפלות. ככל שדיוק הקלט נמוך יותר, כך התפוקה גבוהה יותר - וזו הסיבה שבכל דור NVIDIA מוסיפה פורמט "רזה" יותר.

פורמט קלט ביט הוצג בדור הערה
FP16 16 Volta half precision קלאסי, צבירת FP32
BF16 16 Ampere טווח מעריך רחב כמו FP32, פחות מנטיסה - נוח לאימון
TF32 19 (מאוחסן ב-32) Ampere תחליף שקוף ל-FP32 בכפל מטריצות, כמעט חינם
INT8 8 Turing להסקה (inference) מכומתת
FP8 (E4M3 / E5M2) 8 Hopper שני וריאנטים, טווח מול דיוק
FP4 4 Blackwell הדיוק הרזה ביותר, תפוקת שיא

שימו לב ל-TF32 (TensorFloat-32): זהו פורמט של 19 ביט שמאוחסן במקום של FP32 אך נחתך במנטיסה. הוא מאפשר להריץ קוד שנכתב ל-FP32 על ליבות הטנזור בלי לשנות שורת קוד, ברווח תפוקה גדול ובאובדן דיוק זניח - ולכן הוא ברירת המחדל של cuBLAS בהרבה מסלולים. את ההשלכות של בחירת הדיוק על מודל ה-Roofline (arithmetic intensity מול רוחב פס) נעמיק בפרק הביצועים; כאן חשוב לזכור שכל שורה בטבלה היא זוג "קלט זול, צבירה יקרה", וכל דור מוסיף שורה רזה יותר למטה.

חומרת התמיכה - warpgroup, TMA וזיכרון טנזור

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

קבוצת warps - warpgroup (Hopper). ראינו שהוראת wmma היא שיתופית ברמת warp בודד. Hopper הוסיפה יחידת עבודה גדולה יותר - ה-warpgroup, קבוצה של ארבעה warps רצופים, כלומר 128 threads. ההוראה wgmma.mma_async (warpgroup MMA, אסינכרונית) מבצעת כפל מטריצות בtile גדול הרבה יותר, שמשותף על פני כל 128 הthreads. מדוע זה חשוב? כי tiles גדולים יותר רוֹוים (saturate) את רוחב הפס האריתמטי העצום של ליבות הטנזור בכרטיסי מרכזי-נתונים. הtile הקטן משאיר את המנוע רעב; tile בגודל warpgroup ממלא אותו. בנוסף, ה-wgmma היא אסינכרונית - היא משוגרת וממשיכה לרוץ בזמן שהthreads עושים עבודה אחרת - וכך היא מבטלת את תקורת הסנכרון המפורש בין ארבעת ה-warps. זו לבנת יסוד בקרנלים מודרניים כמו Flash Attention.

מאיץ זיכרון הטנזור - TMA (Tensor Memory Accelerator, Hopper/Blackwell). אם ליבת הטנזור טוחנת tiles בקצב מטורף, מי מביא לה אותם? בעבר, כל העתקה מ-global memory (זיכרון גלובלי) ל-shared memory (זיכרון משותף) עברה דרך האוגרים ודרשה מה-CUDA Cores לחשב כתובות. ה-TMA הוא מנוע העתקה ייעודי ואסינכרוני שמעביר נתונים ישירות מזיכרון ה-GPU אל ה-shared memory / cache L1, תוך עקיפה מלאה של קובץ האוגרים. יש לו שני יתרונות גדולים:

  1. חיסכון במשאבים - הורדת חישוב הכתובות מה-CUDA Cores. ה-TMA מחשב בחומרה את הכתובות של גישות "affine" בכמות גדולה, מהצורה addr = width * base + offset. בכך הוא מפנה אוגרים (פחות אוגרים תפוסים באחזקת אינדקסים - פחות register pressure, פחות פגיעה ב-occupancy) ומפנה את ה-CUDA Cores לחישוב אמיתי במקום לאריתמטיקת כתובות. החיסכון בולט במיוחד בגישות גדולות, בסדר גודל של קילובייטים, למערכים רב-ממדיים.
  2. מודל אסינכרוני של יצרן-צרכן. thread בודד יכול לשגר העתקה גדולה ואז לחזור להצטרף ל-warp שלו ולעשות עבודה אחרת; הthreads מזהים מאוחר יותר, אסינכרונית, שההעתקה הסתיימה. זהו מודל יצרן-צרכן קלאסי (ההעתקה היא היצרן, החישוב הוא הצרכן) שמאפשר לחפוף תנועת זיכרון עם חישוב.

זיכרון טנזור - Tensor Memory (Blackwell, B200). בדור Blackwell נוסף מאגר אחסון ייעודי חדש בתוך ה-SM: זיכרון הטנזור (TMEM), שנועד להחזיק את הקלטים והפלטים של ליבות הטנזור - ובמיוחד את מטריצות הצבירה (accumulators). הרציונל, כפי שה-glossary מנסח: "המצברים נגישים בתדירות גבוהה יותר במהלך כפל-המטריצות מאשר הtiles, ולכן הם נהנים יותר מחומרה ייעודית" - חיווט קצר ופשוט יותר. בהוראה tcgen05.mma של Blackwell (המחשבת D += A @ B), חלוקת האופרנדים נוקשה: המצבר D חייב לשבת בזיכרון הטנזור; A יכול לשבת בזיכרון הטנזור או ב-shared memory; B חייב לשבת ב-shared memory (אסור לו בזיכרון הטנזור).

המלכוד הגדול - ה-TMA אינו מאיץ את זיכרון הטנזור

עכשיו, שימו לב היטב, כי כאן מסתתר אחד המלכודים הבזויים ביותר בכל התחום, ו-NVIDIA עצמה הטמינה אותו בבחירת השמות. השמות Tensor Memory Accelerator (TMA) ו-Tensor Memory (TMEM) דומים עד כדי בלבול, אבל:

ה-Tensor Memory Accelerator (TMA) אינו מאיץ את פעולות זיכרון הטנזור (Tensor Memory). הם שני דברים שונים לחלוטין. ה-TMA טוען נתונים אל תוך ה-shared memory / cache L1 - הוא אינו נוגע בזיכרון הטנזור כלל. הנתונים מגיעים אל זיכרון הטנזור בדרך אחרת לגמרי - רק דרך פעולות של ליבות הטנזור עצמן (או טעינה מ-shared memory ומהאוגרים ברמת warpgroup), לא דרך ה-TMA.

בואו נשרטט את מסלול הנתונים המלא כדי לקבע את ההבחנה:

   GPU memory (HBM3, 80 GiB)
            |
            |  <-- the TMA copies here (asynchronous, bypasses registers)
            v
   shared memory / L1 cache  (256 KiB per SM on H100)
            |
            |  <-- via Tensor Cores / MMA instructions only
            v
   Tensor Memory - TMEM  (Blackwell only; mainly holds accumulators)
            |
            |  <-- output goes on to further processing
            v
   activation functions / non-linearities in the neural network

   note: the TMA "skips" only the top arrow.
   it never writes directly to TMEM. the name is deliberately misleading.

הדרך לזכור: ה-TMA הוא מנוע העתקה ("Accelerator" של גישה לזיכרון), ופועלו הוא להביא נתונים אל ה-shared memory. זיכרון הטנזור הוא מקום אחסון בתוך ה-SM, ורק ליבות הטנזור ממלאות ומרוקנות אותו. אם תבלבלו ותחשבו ש"ה-TMA מזרים ישירות לזיכרון הטנזור", תבנו מודל מנטלי שגוי של כל מסלול הנתונים. את המסלול הזה בשלמותו, ואיך CUTLASS ו-CuTe מנצלות אותו, נראה בפרקים 6 עד 8.

סיכום

  • ליבת טנזור (Tensor Core) היא ליבת GPU שפועלת על מטריצות שלמות בכל הוראה ומבצעת כפל-וצבירה D = A*B + C; היא בסגנון CISC (הוראה אחת, המון אריתמטיקה) בניגוד ל-CUDA Core בסגנון RISC, ומספקת כלל-אצבע של פי ~100 בתפוקת ה-FLOPS.
  • ליבות הטנזור גדולות פיזית ומעטות בהרבה מ-CUDA Cores: ב-H100 יש רק ארבע ל-SM, אחת לכל warp scheduler; הן הוצגו לראשונה בדור Volta (V100).
  • ההוראה HMMA16.16816.F32 מקבלת קלט FP16, צוברת ב-FP32, tile m=16 n=8 k=16, ולכן מבצעת 16*8*16 = 2,048 MAC בהוראה אחת, שהם 2,048/32 = 64 MAC לכל thread; היעילות באה מפרישת עלות שליפת ההוראה על פני המון חישוב.
  • הוראת ליבת טנזור מבוצעת שיתופית על ידי warp שלם - אף thread בודד אינו מחשב מטריצה שלמה, הtile פרוס על אוגרי 32 הthreads; מכאן השם warp-level (wmma).
  • שרשרת התכנות היא בת שלוש רמות: wmma::mma_sync ב-CUDA C++, הוראת wmma.mma.sync... אחת ב-PTX, ופיצול לכמה הוראות HMMA קטנות ב-SASS (למשל ארבע HMMA.1688 שמכסות m16n16k16); המיקרו-ארכיטקטורה הפנימית קניינית, ככל הנראה מערך סיסטולי, אך אין קונסנזוס.
  • אין לכתוב הוראות ליבת טנזור ביד; משתמשים ב-cuBLAS, CUTLASS או CuTe DSL (פרק 6), וזוכרים ש-cuBLAS היא column-major.
  • ב-Hopper נוספה ה-warpgroup (ארבעה warps, 128 threads) עבור wgmma.mma_async - tiles גדולים יותר שרווים את ליבות הטנזור ורצים אסינכרונית.
  • ה-TMA (Hopper/Blackwell) הוא מנוע העתקה אסינכרוני מזיכרון ה-GPU אל ה-shared memory תוך עקיפת האוגרים, שמחשב כתובות affine בחומרה ובכך מפנה אוגרים ו-CUDA Cores; זיכרון הטנזור (Blackwell, B200) הוא מאגר על-SM לאופרנדים של ליבות הטנזור, בעיקר למצברים.
  • המלכוד הקריטי: ה-Tensor Memory Accelerator אינו מאיץ את ה-Tensor Memory - הוא טוען ל-shared memory / L1, לא ל-TMEM; הנתונים מגיעים ל-TMEM רק דרך ליבות הטנזור.