1.1 מבנה ה GPU GPC, TPC ו SM הרצאה

בפרק ההקדמה בנינו את התמונה הרעיונית: ראינו בשיעור 0.1 שה-GPU איננו "מהיר יותר" מהמעבד אלא מוכוון-תפוקה (throughput) במקום latency, ושהכוח שלו נובע ממערך גדול של יחידות פשוטות שנקראות SM - מעבד רב-זרמי (Streaming Multiprocessor). עכשיו אנחנו פותחים את המכסה ומסתכלים על החומרה עצמה. בשיעור הזה נבנה את היררכיית ההכלה הפיזית של GPU של NVIDIA - איך מרכיב מכיל את חברו, מהשבב כולו ועד לליבה הבודדת - ונבין למה המבנה הזה נראה כפי שהוא נראה. נתחיל מהרעיון ההיסטורי שהוליד את כל הארכיטקטורה (החזון של Lindholm ושותפיו מ-2008), נרד דרך ה-GPC וה-TPC אל ה-SM, נספור כמה יש מכל דבר ב-H100 שלנו, ולבסוף נראה תופעה מרתקת: איך רבדים שהיו במשך שנים "חומרה בלבד" - ה-GPC וה-TPC - נעשו לאחרונה גלויים למתכנת עצמו בדורות Hopper ו-Blackwell. בשיעור 1.2 נצלול פנימה אל תוך ה-SM ונפרק אותו לרכיביו; כאן אנחנו בונים את השלד שמחזיק את כל ה-SM-ים יחד.

מארכיטקטורת pipeline קבועה לארכיטקטורה מאוחדת - CUDA

לפני שנספור רכיבים, צריך להבין רעיון אחד שמסביר את כל השאר. ראשי התיבות CUDA מייצגים Compute Unified Device Architecture - ארכיטקטורת הdevice המאוחדת לחישוב - וכאן המילה החשובה היא Unified (מאוחדת). כדי להבין מה אוחד, צריך לראות מה היה קודם.

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

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

אחרי CUDA - יחידות אחידות. החזון, שנוסח במאמר הלבן המכונן של Lindholm ואחרים, 2008 (מאמר ארכיטקטורת Tesla, המקור המקורי לחלק גדול מהתרשימים והניסוחים שבתיעוד של NVIDIA עד היום), היה פשוט להפליא: במקום יחידות מתמחות שונות, נבנה יחידות חומרה אחידות לחלוטין, שכל אחת מהן מסוגלת לקשת רחבה של חישובים. היחידות האחידות האלה הן ה-SM-ים. אותה יחידה יכולה בבוקר לחשב פיקסלים ובערב להריץ שכבה של רשת נוירונים.

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

שימו לב לעומס המשמעות של המילה CUDA, כי היא רודפת את הקורס כולו: אותם ארבע אותיות משמשות לשלושה דברים שונים - ארכיטקטורת החומרה (הנושא של השיעור הזה), מודל התכנות (thread, block, grid - פרק 2), ופלטפורמת התוכנה שמרחיבה את C (ה-CUDA Toolkit). כשנאמר "CUDA" בשיעור הזה, הכוונה תמיד לארכיטקטורת החומרה.

היררכיית ההכלה - מ-GPU ועד ליבה

ה-SM איננו צף לבדו על השבב. NVIDIA מארגנת את מאות ה-SM-ים בשרשרת הכלה בת כמה רבדים, מהגדול לקטן. הנה השרשרת המלאה, ואחריה נסביר כל רובד:

GPU (the entire device, e.g. H100 SXM)
│
├── GPC  - GPU Processing Cluster  (Graphics/GPU Processing Cluster)
│    │     includes one raster engine (rasterization engine, a leftover from the graphics past)
│    │
│    ├── TPC  - Texture Processing Cluster  (Texture Processing Cluster)
│    │    │     = a pair of physically adjacent SMs
│    │    ├── SM 0
│    │    └── SM 1
│    │
│    ├── TPC ...
│    └── (up to 9 TPCs per GPC in GH100)
│
├── GPC ...
└── (8 GPCs in the full GH100)

ובתוך כל SM בודד יושבות יחידות הביצוע עצמן, שאותן נפרק לעומק בשיעור 1.2:

SM (a single streaming multiprocessor)
├── 4 × Warp Scheduler   (warp scheduler, one per SM quarter)
├── 128 × CUDA Core       (CUDA core, a scalar FP32 unit)
│    64 × INT32   ·   64 × FP64   ·   4 × Tensor Core (tensor core)
├── SFU / LSU              (special function and load/store units)
├── Register File          (register file, 256 KB per SM)
└── L1 / shared memory     (shared memory + L1 cache, 256 KiB per SM)

שלוש הבחנות שחשוב לקבע כבר עכשיו, כי הן מקור נפוץ לבלבול:

• הליבה איננה המקבילה של ליבת מעבד. ה-CUDA Core היא pipeline חישוב בודד, לא "מעבד". המקבילה האמיתית של ליבת CPU - יחידה עם מתזמנים, אוגרים וcache משלה - היא ה-SM כולו. הזכרנו זאת בשיעור 0.1, וההיררכיה למעלה מראה זאת מבנית: ה-SM הוא הרובד שבו יושבים המתזמנים והזיכרון.
• ה-GPC וה-TPC הם רבדים ארגוניים של החומרה, ובמשך רוב ההיסטוריה של CUDA הם היו בלתי נראים למתכנת - הם לא הופיעו לא במודל התכנות ולא בהיררכיית הזיכרון. בהמשך השיעור נראה איך זה השתנה.
• המספרים המדויקים משתנים משבב לשבב. כמה GPC-ים? כמה TPC-ים ל-GPC? כמה SM-ים בסך הכול? אלה אינם קבועי טבע - הם נקבעים לכל דגם, וקוראים אותם מה-whitepaper של הארכיטקטורה או ישירות מהכרטיס עם deviceQuery. תכף נעשה בדיוק את זה ל-H100.

אשכול עיבוד ה-GPU - GPC

ה-GPC הוא הרובד הגבוה ביותר מתחת ל-GPU עצמו. ראשי התיבות מקורם ב-Graphics Processing Cluster (אשכול עיבוד גרפי), אבל התיעוד העדכני של NVIDIA (למשל ה-CUDA C++ Programming Guide) פותח אותם היום כ-GPU Processing Cluster - שינוי שמסגיר את המעבר של הכרטיס מעולם הגרפיקה אל עולם החישוב הכללי.

מה יש בתוך GPC? שני דברים:

1. אוסף של TPC-ים (ובתוכם ה-SM-ים) - זהו לב החישוב.
2. מנוע רסטריזציה אחד (raster engine) - היחידה שממירה משולשים לפיקסלים. זהו הרכיב הגרפי ה"היסטורי" שנתן ל-GPC את שמו, ובעומסי חישוב טהורים (רשתות נוירונים, סימולציות) הוא כמעט לא רלוונטי.

מבנית, אם כן, ה-GPC הוא "מיני-GPU": צביר של יחידות חישוב עם מנוע רסטריזציה משלו. שבב שלם הוא פשוט כמה GPC-ים זה לצד זה. ב-GH100 (השבב של ה-H100) יש 8 GPC-ים.

עד לפני דור אחד, זה היה כל הסיפור: ה-GPC היה פרט ארגוני של החומרה, ולמתכנת לא היתה שום דרך "לפנות" אליו או להתייחס אליו בקוד. זה השתנה ב-Hopper, ונחזור לזה בסעיף על thread block clusters - שם ה-GPC יהפוך לראשונה ליעד תזמון שהמתכנת יכול לבקש במפורש.

זוג ה-SM הצמוד - TPC

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

המילה החשובה בהגדרה היא צמודים (adjacent). זו לא צמידות מקרית או נוחות תרשימית - הקִרבה הפיזית בין שני ה-SM-ים היא תכונה מהותית. שני SM-ים שיושבים זה ליד זה על השבב יכולים לחלוק מסלולי חיווט קצרים ביניהם, וזה בדיוק מה שיאפשר, בדור Blackwell, פעולה שבה שני ה-SM-ים של אותו TPC משתפים פעולה על אותו חישוב. בלי הצמידות הפיזית, שיתוף פעולה כזה בעלות סבירה לא היה אפשרי. נחזור לזה בסעיף האחרון.

כמה SM-ים ב-TPC? תמיד 2 - זוג, בלי שונות בין הדורות. כמה TPC-ים ב-GPC? זה כן משתנה: ב-GH100 יש עד 9 TPC-ים לכל GPC. שימו לב שההגדרה "זוג SM" מסבירה מיד למה מספר ה-SM-ים בשבב הוא כמעט תמיד זוגי, ולמה הוא נוטה להתחלק יפה: הוא נבנה מיחידות של 2.

היסטורית, השם "עיבוד טקסטורות" מסגיר גם הוא את העבר הגרפי (טקסטורות הן התמונות שנמרחות על משטחים בתלת-ממד). אבל בדיוק כמו ה-GPC, ה-TPC היה במשך שנים מבנה חומרה בלבד, ללא כל מיפוי למודל התכנות - לא לזיכרון ולא לthreads. המתכנת לא ידע ולא היה צריך לדעת שקיים דבר כזה. גם זה השתנה, ב-Blackwell, וזו תהיה הפעם הראשונה שה-TPC ניתן למיעון בקוד.

ה-SM - היחידה המאוחדת בפעולה

הגענו לרובד שבו קורה החישוב האמיתי: ה-SM. הוא היחידה האחידה שעליה דיבר החזון של Lindholm - העותק שמשוכפל מאות פעמים על השבב. כאן רק נציב אותו במקומו בהיררכיה ונזכיר את המספרים המרכזיים; הפירוק המלא של רכיביו (מתזמני ה-warp, קובץ האוגרים, יחידות הביצוע) שמור לשיעור 1.2.

ה-SM נמצא בתוך TPC (זוג), שנמצא בתוך GPC, שנמצא בתוך ה-GPU. הוא מכיל בתוכו את יחידות הביצוע: ב-H100, כל SM כולל 4 מתזמני warp, 128 יחידות CUDA Core מסוג FP32 (וכמחצית מזה יחידות INT32 ו-FP64), 4 Tensor Cores (אחד לכל מתזמן), קובץ אוגרים של 256 KB, ו-256 KiB של זיכרון L1/shared משותף. הסיבה שקוראים לו "מעבד רב-זרמי" היא שהוא מריץ בו-זמנית זרמים רבים של threads - עד 64 warps (2,048 threads) תושבים בכל רגע - ומחליף ביניהם בעלות של מחזור שעון אחד כדי להסתיר latency, בדיוק כפי שראינו בשיעור 0.1.

מבחינת ההיררכיה של השיעור הזה, מספיק לזכור: ה-SM הוא היחידה שהמתכנת חושב עליה כ"ליבה" האמיתית, והוא היעד שאליו מתזמן החומרה משבץ block שלם. בדיוק כפי ש-block משובץ ל-SM יחיד, נראה מיד שאשכול של block-ים (cluster) משובץ ל-GPC יחיד - אנלוגיה שתעשה את כל הרובד של ה-GPC לפתע רלוונטי לתכנות.

כמה יש מכל דבר ב-H100 - קריאה מ-whitepaper ומ-deviceQuery

עכשיו נמלא את המספרים המדויקים לכרטיס הייחוס שלנו. ה-H100 SXM כולל 132 SM-ים. מספר זה איננו מספר שצריך לזכור בעל פה - צריך לדעת מאיפה לקרוא אותו. שני מקורות אמת:

מקור 1 - ה-whitepaper של הארכיטקטורה. המסמך "NVIDIA H100 Tensor Core GPU Architecture" מפרט את מבנה שבב ה-GH100. השבב המלא (full GH100) בנוי כך:

Full GH100:
  8 GPCs
  × 9 TPCs per GPC        = 72 TPCs
  × 2 SMs per TPC         = 144 SMs  (in the full chip)

H100 SXM in practice:  132 SMs enabled  (66 TPCs)

למה 132 ולא 144? כי בייצור שבבים ענקיים תמיד יש פגמים נקודתיים. NVIDIA מכבה חלק מה-SM-ים הפגומים ומוכרת את השבב עם 132 SM-ים מתפקדים במקום 144. זו הסיבה שכמעט תמיד מספר ה-SM-ים בכרטיס מסחרי קטן במקצת מהמקסימום התאורטי של השבב - עובדה שנקראת מה-whitepaper, לא מנחשים אותה.

מקור 2 - הכרטיס עצמו, דרך deviceQuery. תוכנית הדוגמה deviceQuery מ-CUDA Samples שואלת את ה-runtime ומדפיסה, בין השאר:

Device 0: "NVIDIA H100 80GB HBM3"
  CUDA Capability Major/Minor version number:    9.0
  (132) Multiprocessors, (128) CUDA Cores/MP:    16896 CUDA Cores
  Total amount of global memory:                 81559 MBytes

השורה (132) Multiprocessors היא בדיוק מספר ה-SM-ים. שימו לב שהיא גם מכפילה: 132 × 128 = 16,896 יחידות CUDA Core מסוג FP32 - אותו מספר הthreads המקביליים-באמת שחישבנו בשיעור 0.1. את המספר הזה אפשר לקרוא גם בתוך CUDA C++ עם השדה multiProcessorCount:

#include <cstdio>
#include <cuda_runtime.h>

#define CUDA_CHECK(call)                                                      \
    do {                                                                      \
        cudaError_t err_ = (call);                                            \
        if (err_ != cudaSuccess) {                                            \
            fprintf(stderr, "CUDA error at %s:%d\n  '%s'\n  -> %s\n",         \
                    __FILE__, __LINE__, #call, cudaGetErrorString(err_));     \
            exit(EXIT_FAILURE);                                               \
        }                                                                     \
    } while (0)

int main(void) {
    cudaDeviceProp p;
    CUDA_CHECK(cudaGetDeviceProperties(&p, 0));
    printf("%s  (compute capability %d.%d)\n", p.name, p.major, p.minor);
    printf("SM count (multiProcessorCount): %d\n", p.multiProcessorCount);
    return 0;
}

הcompilation והרצה (sm_90a היא הארכיטקטורה של Hopper; החליפו לפי הכרטיס שלכם):

nvcc -O2 -arch=sm_90a -o smcount smcount.cu
./smcount

הפלט על H100:

NVIDIA H100 80GB HBM3  (compute capability 9.0)
SM count (multiProcessorCount): 132

ומה עם GPC ו-TPC? הנה נקודה חשובה: ה-runtime של CUDA אינו חושף ישירות את מספר ה-GPC-ים או ה-TPC-ים. אין שדה gpcCount ב-cudaDeviceProp. את החלוקה הזו קוראים מה-whitepaper בלבד. מה ש-deviceQuery כן נותן זה multiProcessorCount, ומתוכו, בעזרת הידע ש-TPC = 2 SM, אפשר להסיק שמספר ה-TPC-ים המופעלים הוא 132 / 2 = 66.

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

כרטיס	ארכיטקטורה	compute capability	SM-ים (מופעלים)	SM-ים בשבב מלא	הערה על המבנה
A100 SXM	Ampere (GA100)	8.0	108	128	8 GPC × עד 8 TPC × 2 SM
H100 SXM	Hopper (GH100)	9.0	132	144	8 GPC × עד 9 TPC × 2 SM
B200	Blackwell (GB100)	10.0	על סדר גודל של ~160 (שני dies)	תלוי SKU	שני dies הפועלים כ-GPU אחד

שימו לב ל-B200: הוא מיוצר משני dies נפרדים המחוברים בקישור מהיר ופועלים כ-GPU לוגי אחד. מכיוון שהספירה תלויה ב-SKU ובכמה SM-ים כובו, את המספר המדויק קוראים מה-whitepaper של Blackwell ומ-deviceQuery - וזה בדיוק הלקח החוזר: לא זוכרים מספרים, יודעים מאיפה לשלוף אותם. הקבוע היחיד לאורך כל השורות בטבלה הוא ש-TPC הוא תמיד זוג SM ושהמבנה הוא תמיד GPU → GPC → TPC → SM.

כשה-GPC נעשה גלוי לתוכנה - thread block clusters ב-Hopper

עד כה תיארנו את ה-GPC וה-TPC כרבדים של חומרה בלבד. ב-Hopper (compute capability 9.0, ה-H100) זה השתנה לראשונה עבור ה-GPC. NVIDIA הוסיפה רובד חדש להיררכיית הthreads של CUDA: מעל ה-block, אבל מתחת ל-grid, נכנס ה-thread block cluster - אשכול של block-ים.

הרעיון בנוי על אנלוגיה מדויקת:

a single block      →  scheduled onto a single SM
a cluster of blocks →  scheduled onto a single GPC

כלומר, בדיוק כפי שהחומרה מבטיחה שכל הthreads של block מסוים ירוצו יחד על אותו SM (וכך יכולים לחלוק את ה-shared memory שלו), ב-Hopper החומרה מבטיחה שכל ה-block-ים של אותו cluster ירוצו יחד על אותו GPC. וזו בדיוק הסיבה שה-GPC נעשה לפתע רלוונטי למתכנת: הוא הפך ליעד תזמון שאפשר לבקש.

מה מרוויחים מזה? זיכרון משותף מבוזר (distributed shared memory). מכיוון שכל ה-block-ים בקלאסטר יושבים על SM-ים באותו GPC, החומרה מאפשרת ל-block אחד לגשת ישירות ל-shared memory של block אחר באותו cluster - רובד חדש בהיררכיית הזיכרון, בין ה-shared memory המקומי של SM בודד לבין ה-global memory האיטי. הנה שלד קוד שמדגים את שני החלקים - הכרזת ממדי הקלאסטר וגישה מבוזרת דרך cooperative groups:

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

// A cluster of 2 blocks; the hardware will schedule both onto the same GPC
__global__ void __cluster_dims__(2, 1, 1) clusterKernel(int* out) {
    cg::cluster_group cluster = cg::this_cluster();
    __shared__ int smem[256];

    unsigned int rank = cluster.block_rank();           // 0 or 1 within the cluster
    smem[threadIdx.x] = rank;                            // each block writes to its own memory
    cluster.sync();                                      // barrier across the entire cluster

    // access to the shared memory of the neighboring block in the same GPC (distributed shared memory):
    unsigned int peer = (rank + 1) % cluster.num_blocks();
    int* remote = cluster.map_shared_rank(smem, peer);
    out[blockIdx.x * blockDim.x + threadIdx.x] = remote[threadIdx.x];
}

הlaunch נעשית עם -arch=sm_90a (הסיומת a מפעילה את התכונות הייחודיות ל-Hopper), וגדלי הקלאסטר נתמכים גם דרך cudaLaunchKernelEx עם התכונה cudaLaunchAttributeClusterDimension כשרוצים לקבוע אותם בזמן ריצה. נעמיק ב-API עצמו בפרק על הזיכרון; כאן הנקודה המבנית: זו הפעם הראשונה שה-GPC "עלה" מרובד החומרה הבלתי-נראה אל תוך מודל התכנות. מי שכתב CUDA לפני Hopper מעולם לא כתב מילה על GPC; מ-Hopper ואילך, כשמבקשים cluster, מבקשים בעצם "שבץ את הכול על GPC אחד".

כשה-TPC נעשה גלוי לתוכנה - tcgen05 ו-.cta_group ב-Blackwell

הדור הבא, Blackwell (Tensor Cores מהדור החמישי), עשה את אותו הדבר לרובד אחד למטה: הוא הפך את ה-TPC לגלוי בקוד, וכך נסגר המעגל - כל רובד בהיררכיה, מלמעלה למטה, קיבל בסופו של דבר ביטוי תוכנתי.

הזכרנו שה-TPC הוא זוג SM צמודים, ושהצמידות הפיזית נועדה לאפשר לשני ה-SM-ים לשתף פעולה. ב-Blackwell זה מתממש דרך משפחת פקודות ה-PTX החדשה tcgen05 (פקודות ה-Tensor Core של הדור החמישי). הפקודות האלה קיבלו שדה חדש בשם .cta_group שקובע האם הפעולה מכוונת ל-SM בודד או לזוג ה-SM-ים של ה-TPC:

מְאַפְיֵן ב-PTX	היעד הפיזי	וריאנט ה-SASS
.cta_group::1	SM בודד	1SM
.cta_group::2	זוג ה-SM-ים של ה-TPC	2SM

הדוגמה הבולטת היא פקודת ה-MMA (matrix multiply-accumulate) - כפל-והצטברות של מטריצות, לב-לבו של ה-Tensor Core. באותו קוד PTX, המתכנת (או הcompiler של הספרייה) בוחר האם להנפיק את הכפל ל-SM בודד או לגייס את שני ה-SM-ים של ה-TPC לעבוד יחד על אותה מטריצה. באופן סכמטי (האופרנדים הושמטו לשם הבהירות):

// matrix multiply on a single SM:
tcgen05.mma.cta_group::1.kind::f16   [tmem_d], desc_a, desc_b, ... ;

// the same multiply, recruiting the TPC's SM pair (2SM cooperation):
tcgen05.mma.cta_group::2.kind::f16   [tmem_d], desc_a, desc_b, ... ;

ההבדל בין .cta_group::1 ל-.cta_group::2 מתורגם על ידי הcompiler לשתי גרסאות SASS שונות (1SM מול 2SM) - כלומר, בחירת המְאַפְיֵן היא החלטת קומפילציה ששולטת בהתנהגות בזמן ריצה. כאן הצמידות הפיזית שדיברנו עליה משתלמת: מכיוון ששני ה-SM-ים של ה-TPC שכנים על הסיליקון, הם יכולים לחלק ביניהם את המטריצה, לחלוק אופרנדים דרך המסלולים הקצרים, ולהכפיל בפועל את המשאבים הזמינים לפעולת MMA בודדת - ללא העלות שהיתה כרוכה בשיתוף בין SM-ים רחוקים.

וכך, מ-Blackwell ואילך, ה-TPC ניתן למיעון בקוד לראשונה. שימו לב מה קרה על פני הדורות, כי זה הסיפור הגדול של השיעור:

Until Hopper:  GPC, TPC  =  hardware only, invisible to software
Hopper (9.0):  GPC        →  became visible via thread block clusters (+ distributed shared memory)
Blackwell:     TPC        →  became visible via tcgen05 .cta_group  (1SM / 2SM)

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

סיכום

• ה-CUDA כארכיטקטורת חומרה מבוססת על חזון Lindholm ואחרים 2008: החלפת ה-pipeline הקבוע של יחידות מתמחות הטרוגניות ביחידות אחידות (SM-ים) שכל אחת מסוגלת לקשת רחבה של חישובים - פישוט בו-זמני לתוכנה (אין מיפוי לpipeline קבוע) ולחומרה (אין ניחוש יחסי עומס בין שלבים).
• היררכיית ההכלה הפיזית היא GPU → GPC → TPC → SM → (CUDA Cores, Tensor Cores, ...); ה-GPC כולל גם raster engine, וה-TPC הוא תמיד זוג SM.
• ה-TPC מוגדר כזוג SM צמודים פיזית, והצמידות איננה מקרית - היא מה שמאפשר בהמשך את שיתוף הפעולה בין שני ה-SM-ים (וריאנט 2SM ב-Blackwell).
• ה-H100 SXM כולל 132 SM-ים מתוך 144 בשבב GH100 המלא (8 GPC × עד 9 TPC × 2 SM); ההפרש נובע מ-SM-ים פגומים שכובו בייצור.
• את מספר ה-SM-ים קוראים, לא זוכרים: מה-whitepaper של הארכיטקטורה, מ-deviceQuery ((132) Multiprocessors), או מהשדה multiProcessorCount ב-CUDA C++; מספר ה-GPC/TPC אינו נחשף על ידי ה-runtime ונקרא מה-whitepaper בלבד.
• הספירות משתנות בין דגמים: A100 עם 108 SM, H100 עם 132 SM, ו-B200 (שני dies) על סדר גודל של כ-160 SM; הקבוע היחיד הוא המבנה GPU → GPC → TPC(זוג SM) → SM.
• ב-Hopper (compute capability 9.0) ה-GPC נעשה גלוי לתוכנה לראשונה דרך thread block clusters: אשכול של block-ים משובץ ל-GPC יחיד (כפי ש-block משובץ ל-SM יחיד), מה שמכניס distributed shared memory כרובד חדש בהיררכיית הזיכרון.
• ב-Blackwell ה-TPC נעשה גלוי לתוכנה דרך שדה ה-.cta_group בפקודות ה-PTX מסוג tcgen05: .cta_group::1 מכוון ל-SM בודד (SASS מסוג 1SM) ו-.cta_group::2 מכוון לזוג ה-SM-ים של ה-TPC (2SM), למשל בפקודת MMA - הפעם הראשונה שה-TPC ניתן למיעון בקוד.