8.5 Register pressure ואיזון משאבים הרצאה

לאורך פרק 8 בנינו את מדע ה-occupancy (תפוסה). בשיעור 8.1 חישבנו occupancy תיאורטי מתצורת הlaunch ומגבלות ה-SM, וראינו שהמשאבים הסופיים של ה-SM - קובץ האוגרים (register file) וה-shared memory - הם שקובעים כמה warps תושבים נכנסים. בשיעור 1.6 הכרנו את קובץ האוגרים כחומרה: זיכרון SRAM מהיר פי עשרה מה-L1, שמחזיק את האופרנדים בזמן שהליבות מחשבות עליהם, ושמחולק בין כל ה-threads התושבים על ה-SM. בשיעור הזה נסגור את הפרק ואת המעגל: נראה מה קורה כשקובץ האוגרים עצמו הופך לצוואר-הבקבוק - התופעה שנקראת בשם הציורי register pressure (לחץ אוגרים). נבין מדוע כל אוגר נוסף שכל thread צורך גובה מחיר ב-occupancy, מדוע דליפה (spilling) של אוגרים לזיכרון היא אסון שקט, ובעיקר - איך מודדים ושולטים בכל זה עם nvcc -Xptxas -v, עם __launch_bounds__ ועם -maxrregcount. התובנה המרכזית, שתלווה אותנו עד הסוף, היא שאין כאן "יותר זה תמיד יותר טוב": יותר אוגרים ל-thread מאיצים thread בודד אבל מורידים occupancy, ונקודת האיזון נמצאת רק במדידה, לא בנוסחה.

לחץ אוגרים - register pressure

ה-glossary מגדיר את המונח במשפט אחד קולע: register pressure הוא ביטוי ציורי לכך שקובץ האוגרים הוא צוואר-הבקבוק. זהו לא באג ולא שגיאה - זו מתיחות מובנית בארכיטקטורה. מצד אחד, כל thread רוצה כמה שיותר אוגרים: אוגר הוא האחסון המהיר ביותר בהיררכיה, וככל שיש לו יותר אוגרים כך הוא יכול להחזיק יותר ערכי-ביניים "חמים" בלי לחשב אותם מחדש ובלי לגשת לזיכרון. מצד שני, קובץ האוגרים של ה-SM הוא סופי, והוא מחולק בין כל ה-threads התושבים בו-זמנית. כל אוגר שאתם מבזבזים על thread אחד מוכפל במספר ה-threads התושבים, ומכרסם בתקציב שקובע כמה warps בכלל נכנסים.

שרשרת הסיבתיות, שראינו את קצה שלה כבר ב-1.6, היא הלב של השיעור:

   more registers per thread
        │
        ▼
   more registers per warp  ->  fewer warps fit in the register file
        │
        ▼
   fewer resident warps  ->  lower occupancy
        │
        ▼
   fewer eligible warps to issue  ->  weaker latency hiding
        │
        ▼
   the warp scheduler runs out of work  ->  execution units sit idle

זה בדיוק ההיפוך של חוק Little משיעור 7.4: שם הראינו שצריך מספיק warps כדי להסתיר את ה-latency; register pressure הוא הכוח שמונע מכם להגיע למספר הזה. כשכל thread תאב-אוגרים, פשוט לא נכנסים מספיק warps כדי להזין את ה-scheduler, וה-latency שהיה אמור להיות מוסתר מתגלה כזמן-בטלה גלוי.

אוגרים וירטואליים מול פיזיים - ה-ptxas ממפה ביניהם

כדי להבין מאיפה נקבע מספר האוגרים ל-thread, צריך להיזכר בשני המישורים של הקוד משיעורים 4.x. ב-PTX (הרצת שרשורים מקבילית, ה-IR הווירטואלי), האוגרים הם וירטואליים ובלתי-מוגבלים. כשראיתם בשיעור 4.2 הצהרה כמו .reg .f32 %f<7>, המספר 7 הוא רק כמה אוגרים וירטואליים ה-compiler החליט להשתמש בהם בקטע הזה - ואין שום תקרה עליו. PTX הוא מכונה מופשטת; הוא יכול "לדמיין" אלפי אוגרים.

ב-SASS (אסמבלר זורם, קוד המכונה האמיתי של ה-SM), לעומת זאת, האוגרים הם פיזיים ומוגבלים - הם מתייחסים לתאים ממשיים בקובץ האוגרים של ה-SM. מי שמגשר בין שני העולמות הוא ה-ptxas, הcompiler שמתרגם PTX ל-SASS (הכרנו אותו בשיעור 4.4). ה-ptxas לוקח את האוגרים הווירטואליים הבלתי-מוגבלים וממפה אותם על מספר סופי של אוגרים פיזיים - בדיוק כמו שcompiler על CPU עושה register allocation. וכמו על CPU, כשאין מספיק אוגרים פיזיים, משהו חייב "לדלוף" (על כך בהמשך).

   CUDA C++            PTX (virtual)           SASS (physical)
   ────────            ──────────────          ───────────
   float x = ...;      .reg .f32 %f<7>;        R0..R31  (real registers
   float y = ...;      (unbounded)       ptxas              in the register file)
        │                   │           ────►        │
   nvcc frontend       nvcc / -O          mapping     limit: 255 per thread

שתי מסקנות שנגזרות מכאן, וחשוב להפנים אותן:

1. מספר האוגרים ל-thread נקבע על-ידי ה-SASS, כלומר על-ידי ה-ptxas, ולא על-ידי מה שכתבתם ב-CUDA C++ ישירות. אתם משפיעים עליו בעקיפין - דרך מבנה הקוד ודרך דגלים - אבל ההחלטה הסופית היא של ה-compiler.
2. מכיוון שכל ה-threads של block מתוזמנים לאותו SM, סך האוגרים שה-block צורך הוא מספר-האוגרים-ל-thread כפול מספר ה-threads-ל-block. לכן register pressure תלוי גם בקוד (SASS) וגם בתצורת הlaunch (threads/block). אותו kernel בדיוק, עם block של 128 threads מול block של 512 threads, יוצר לחץ שונה לגמרי על ה-SM.

קובץ האוגרים של ה-H100 - החשבון

עכשיו נכניס מספרים אמיתיים. שיעור 1.6 השאיר לנו מכוון את הקווים הכלליים; הנה הקבועים הקונקרטיים של ה-H100 (compute capability 9.0):

• קובץ האוגרים: 65,536 אוגרים של 32 סיביות לכל SM. זהו התקציב הקשיח שמתחלק בין כל ה-threads התושבים.
• מגבלת חומרה ל-thread: 255 אוגרים. אף thread לא יכול להשתמש ביותר, גם אם ירצה - ה-ISA מקצה 8 סיביות למספר האוגר.
• מקסימום 64 warps (2,048 threads) תושבים ל-SM - התקרה שהכרנו ב-8.1.
• הקצאה בגרנולריות של warp, ביחידות של 256 אוגרים. ה-SM לא מקצה אוגרים ל-thread בודד אלא ל-warp שלם (32 threads), ומעגל את הכמות כלפי מעלה לכפולה של 256.

מכאן נגזר החשבון המרכזי של השיעור. אם thread צורך R אוגרים, warp צורך 32 x R אוגרים, מעוגל כלפי מעלה לכפולה של 256. מספר ה-warps שנכנסים הוא 65536 / (אוגרים-ל-warp), חסום ב-64. נציב כמה ערכים:

אוגרים ל-thread (R)	אוגרים ל-warp (מעוגל ל-256)	warps מקסימלי	occupancy
32	1,024	64	100%
40	1,280	51	80%
48	1,536	42	66%
64	2,048	32	50%
96	3,072	21	33%
128	4,096	16	25%
255	8,192	8	12.5%

הטבלה הזו היא כל הסיפור בתמצית. שימו לב לסף המכריע: 32 אוגרים ל-thread הם בדיוק המספר שמאפשר occupancy מלא של 100% ב-H100 (32 x 32 = 1024 אוגרים ל-warp, ו-65536 / 1024 = 64 warps). כל אוגר מעבר ל-32 מתחיל לגזול warps. ב-64 אוגרים - כפליים - ה-occupancy כבר צנח לחצי. וב-255 האוגרים המקסימליים, נכנסים רק 8 warps ל-SM, occupancy של 12.5% בלבד. זו המשמעות המדויקת של register pressure: כל אוגר עולה occupancy, והמחיר לא ליניארי אלא מדרגתי (בגלל העיגול ל-256).

חשוב לזכור: occupancy מוגבל על-ידי המשאב הצר ביותר. האוגרים הם רק אחד משלושה חסמים אפשריים - יחד עם shared memory ל-block ועם התקרה של 64 warps / 32 blocks ל-SM. ה-occupancy בפועל הוא המינימום על פני כל החסמים. בשיעור הזה אנחנו מתמקדים בחסם האוגרים, אבל בכל ניתוח אמיתי בודקים את שלושתם.

מדידה - קריאת הפלט של ptxas -v

אי אפשר לאזן מה שלא מודדים, והכלי הבסיסי כאן חינמי ומיידי: הדגל -v (verbose) של ה-ptxas. מעבירים אותו דרך ה-nvcc באחת משתי צורות שקולות:

nvcc -O2 -arch=sm_90a -Xptxas -v -c kernel.cu -o kernel.o
# fully equivalent:
nvcc -O2 -arch=sm_90a --ptxas-options=-v -c kernel.cu -o kernel.o

הפלט מודפס ל-stderr, שורה לכל kernel, ונראה כך עבור kernel "רגיל" בלי בעיות:

ptxas info    : Compiling entry function '_Z5lightPKfPfi' for 'sm_90a'
ptxas info    : Function properties for _Z5lightPKfPfi
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 16 registers, 372 bytes cmem[0]

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

• Used N registers - מספר האוגרים ל-thread (N). זה המספר שמזין ישירות את טבלת ה-occupancy מהסעיף הקודם. כאן 16 אוגרים -> occupancy מלא בקלות.
• smem (אם מופיע, למשל 8192 bytes smem) - כמות ה-shared memory ל-block. זהו החסם השני על ה-occupancy.
• spill stores / spill loads - כמות הדליפה (בבייטים). בפלט למעלה שניהם 0 - מצוין. כשהם גדולים מ-0, יש לכם צרה, וזה נושא הסעיף הבא.

(ה-cmem[k] הוא constant memory - בעיקר הארגומנטים של ה-kernel ו-__constant__; הוא כמעט אף פעם לא צוואר-בקבוק, אז לרוב מתעלמים ממנו.) הרגל הזהב: הריצו -Xptxas -v על כל kernel חם, והסתכלו על שלושת המספרים לפני כל אופטימיזציה אחרת.

דליפה לזיכרון מקומי - register spilling

מה קורה כשה-ptxas צריך יותר מ-255 אוגרים ל-thread (או כשהגבלתם אותו למספר נמוך יותר)? הוא לא יכול להמציא אוגרים פיזיים, אז הוא דולף (spills): הוא מוציא חלק מהערכים מהאוגרים ומאחסן אותם ב-local memory (זיכרון מקומי). וכאן המלכודת הגדולה: השם "local" מטעה. local memory הוא לא זיכרון מהיר על ה-SM - הוא פשוט אזור פרטי-ל-thread בתוך ה-GPU RAM (זיכרון גלובלי, DRAM/HBM). הוא "מקומי" רק במובן הלוגי שהוא פרטי ל-thread; פיזית, גישה אליו יורדת עד לתחתית ההיררכיה שראינו ב-1.6 (עם cache ב-L1/L2, אבל בעיקרון latency של מאות מחזורים).

לכן דליפה היא כה יקרה, ולכן היא כה מסוכנת: היא שקטה. הקוד נשאר נכון לחלוטין; שום דבר לא נשבר; שום אזהרה לא קופצת בזמן ריצה. פשוט, כל ערך שדלף הופך גישת אוגר (בעצם חינם) לגישת זיכרון גלובלי (מאות מחזורים). kernel יכול להיות איטי פי כמה בגלל דליפה, ואתם לא תדעו - אלא אם הרצתם -Xptxas -v. הנה איך נראה הפלט כשיש דליפה:

ptxas info    : Function properties for _Z5heavyPKfPfi
    336 bytes stack frame, 328 bytes spill stores, 412 bytes spill loads
ptxas info    : Used 128 registers, 372 bytes cmem[0]

spill stores ו-spill loads גדולים מ-0 = דגל אדום. ה-stack frame גם הוא סימן: הוא ה-local memory שה-frame של ה-thread תופס. שימו לב לדינמיקה הלא-אינטואיטיבית: אם תגבילו את מספר האוגרים ל-thread נמוך מדי (כדי להעלות occupancy), אתם עלולים דווקא לגרום לדליפה - ואז שילמתם פעמיים: גם occupancy נמוך יותר לא הרווחתם ריאלית, וגם כל thread עכשיו נכנס-ויוצא מהזיכרון. זו בדיוק המתיחות: לא מעט מדי אוגרים ולא יותר מדי.

   too many registers            too few registers
   ────────────────               ────────────────
   low occupancy                  spilling to local memory
   fewer warps to hide latency    memory accesses instead of registers
        │                              │
        └──────► the balance point ◄───────┘
              (found only by measuring)

שליטה בקומפיילר - launch_bounds ו-maxrregcount

יש לנו שני מנופים עיקריים להנחות את ה-ptxas כמה אוגרים להקצות. הראשון, המומלץ, הוא __launch_bounds__ - תכונה (attribute) שמצמידים להגדרת ה-kernel עצמו:

__global__ void __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM)
myKernel(const float* in, float* out, int n) { /* ... */ }

שני הפרמטרים אומרים ל-compiler:

• maxThreadsPerBlock - הבטחה שלכם שלא תשיקו את ה-kernel עם block גדול יותר. זה מאפשר ל-compiler לדעת את גבול ה-block.
• minBlocksPerSM (רשות) - כמה blocks לפחות אתם רוצים שיֵשבו על SM בו-זמנית. ה-compiler מחשב מזה תקרת-אוגרים ומכוון אליה.

הדוגמה הקונקרטית על H100: __launch_bounds__(256, 8) אומר "block של עד 256 threads, לפחות 8 blocks ל-SM". החשבון של ה-compiler: 256 threads x 8 blocks = 2048 threads = 64 warps - occupancy מלא. תקציב האוגרים ל-thread שנגזר: 65536 / 2048 = 32 אוגרים ל-thread. אז ה-ptxas ינסה בכל כוחו לרדת ל-32 אוגרים (בדיוק סף ה-100% מהטבלה), וידלוף רק אם באמת אין ברירה.

המנוף השני, הגס יותר, הוא הדגל -maxrregcount=N (או --maxrregcount N):

nvcc -O2 -arch=sm_90a -maxrregcount=32 -Xptxas -v kernel.cu -o kernel

הוא כופה תקרה קשיחה של N אוגרים ל-thread על כל ה-kernels ביחידת הקומפילציה. זו גם החולשה שלו: הוא גלובלי וחסר-אבחנה, בעוד ש-__launch_bounds__ מכוון לכל kernel בנפרד. לכן ההמלצה: העדיפו __launch_bounds__ לכל kernel, ושמרו את -maxrregcount לניסויי-סריקה מהירים (כמו שנעשה בתרגול, לסרוק ערכי-אוגרים ולמצוא את המיטבי בלי לגעת בקוד).

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

ההקלה מהחומרה - הזזת נתונים מהאוגרים

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

• דור Ampere - העתקה א-סינכרונית (async copy). במקום שכל thread ישאב נתון מהזיכרון הגלובלי דרך אוגר (load-to-register ואז store-to-shared), ה-cp.async מעביר ישירות מ-global ל-shared memory ומדלג על האוגרים. פחות אוגרים תפוסים בהחזקת נתונים-בדרך = פחות register pressure.
• דור Hopper - מאיץ זיכרון הטנזור - TMA (Tensor Memory Accelerator). ראינו אותו בשיעור 1.x: יחידה ייעודית שמעתיקה מערכים רב-ממדיים מ-GPU RAM ל-shared memory, מחשבת את כתובות-הגישה האפיניות (addr = width x base + offset) בחומרה, ועוקפת את קובץ האוגרים לגמרי. שני רווחים ל-register pressure: לא צריך אוגרים להחזקת אינדקסים וכתובות, ולא צריך אוגרים להחזקת הנתונים בדרך.
• דור Blackwell - זיכרון הטנזור (Tensor Memory). מאגר ייעודי על ה-SM שמחזיק את האופרנדים - ובעיקר את הצוברים (accumulators) - של ה-Tensor Cores, במקום להחזיק אותם בקובץ האוגרים הכללי. matmul גדול היה מציף את קובץ האוגרים בצוברים; Blackwell מוציא אותם לזיכרון נפרד, ומשחרר את קובץ האוגרים לשאר העבודה.

הקו המשותף בולט: ככל שיותר תנועת-נתונים יורדת מהאוגרים ומהליבות אל יחידות ייעודיות, כך register pressure יורד וה-occupancy יכול לעלות - בלי לוותר על מהירות ה-thread הבודד. זו בעצם דרך שלישית לפתור את המתיחות, מעבר לשני מנופי-ה-compiler: לא לאזן טוב יותר בין אוגרים ל-occupancy, אלא להוריד את הביקוש לאוגרים מלכתחילה. (SemiAnalysis כתבו מאמר מצוין המנתח בדיוק את הקשרים האלה, וה-glossary מפנה אליו.)

המאזן - למצוא את הנקודה המיטבית

נחבר הכל לכלל-פעולה. register pressure הוא לא בעיה שיש לה "פתרון" חד-משמעי; הוא סחר-חליפין שיש לו נקודה מיטבית, והנקודה שונה מ-kernel ל-kernel. שני הכיוונים:

• יותר אוגרים ל-thread מאיצים thread בודד: פחות חישובים-מחדש, פחות דליפות, יותר ערכים חמים באחסון המהיר ביותר. המחיר: פחות warps, occupancy נמוך, הסתרת latency חלשה יותר.
• פחות אוגרים ל-thread מעלים occupancy: יותר warps להסתרת latency. המחיר: אם ירדתם נמוך מדי - דליפה, וכל thread עכשיו איטי יותר.

והנה הנקודה הכי חשובה, שסוגרת גם את הפרק וגם את שיעור 7.4: המטרה היא throughput מרבי, לא occupancy מרבי. ראינו בחוק Little שצריך רק מספיק warps כדי להסתיר את ה-latency - למשל, שרשרת load אופיינית מוסתרת כבר ב-occupancy צנוע של כ-20%. מעבר לזה, occupancy נוסף לא קונה כלום, ואם הוא בא במחיר של פחות אוגרים ל-thread (ואולי דליפה) - הוא אפילו מזיק. לכן קורה לא-אחת שה-kernel הכי מהיר הוא זה עם occupancy בינוני: מספיק warps כדי להסתיר latency, אבל מספיק אוגרים כדי שכל thread ירוץ בלי דליפות ובלי חישובים מיותרים. אין נוסחה שנותנת את הנקודה הזו. יש רק את הלולאה: לשנות את מספר האוגרים (עם __launch_bounds__ או -maxrregcount), למדוד occupancy ו-runtime, ולחזור - עד שמוצאים את מספר-האוגרים שממקסם throughput. את הלולאה הזו נתרגל בדיוק בתרגול.

סיכום

• ה-register pressure (לחץ אוגרים) הוא המצב שבו קובץ האוגרים הוא צוואר-הבקבוק: כל אוגר שכל thread צורך מוכפל במספר ה-threads התושבים ומכרסם בתקציב שקובע כמה warps נכנסים ל-SM.
• אוגרים ב-PTX הם וירטואליים ובלתי-מוגבלים, אוגרי ה-SASS פיזיים ומוגבלים, וה-ptxas ממפה בין השניים; מספר האוגרים ל-thread נקבע על-ידי ה-SASS יחד עם תצורת הlaunch (threads/block), כי כל ה-threads של block יושבים על אותו SM.
• ל-H100 יש 65,536 אוגרים של 32 סיביות לכל SM, מגבלת חומרה של 255 אוגרים ל-thread, והקצאה בגרנולריות של warp ביחידות של 256; מכאן ש-32 אוגרים ל-thread נותנים occupancy מלא (64 warps), 64 אוגרים נותנים 50%, ו-255 אוגרים נותנים 12.5% בלבד.
• הדגל nvcc -Xptxas -v (או --ptxas-options=-v) מדפיס לכל kernel את מספר האוגרים, את ה-shared memory, ואת בייטי הדליפה (spill stores / spill loads) - שלושת המספרים שקוראים לפני כל אופטימיזציה.
• דליפה (spilling) מוציאה אוגרים ל-local memory, שהוא למעשה אזור פרטי-ל-thread בתוך ה-GPU RAM (זיכרון גלובלי איטי) - לא זיכרון מהיר על ה-SM; הדליפה שקטה (הקוד נכון אבל איטי) ולכן מסוכנת, ומתגלה רק ב--Xptxas -v.
• __launch_bounds__(maxThreadsPerBlock, minBlocksPerSM) מנחה את ה-compiler לכל kernel בנפרד (מומלץ), ו--maxrregcount=N כופה תקרה גלובלית קשיחה; שניהם מחליפים דליפה ב-occupancy, ולהוריד את התקרה נמוך מדי דווקא גורם לדליפה.
• תכונות חומרה שמקלות על register pressure על-ידי הזזת נתונים מהאוגרים: העתקה א-סינכרונית ב-Ampere, ה-TMA ב-Hopper (עוקף את קובץ האוגרים ומחשב כתובות בחומרה), וזיכרון הטנזור ב-Blackwell (מחזיק צוברים מחוץ לקובץ האוגרים).
• האיזון הוא סחר-חליפין: יותר אוגרים מאיצים thread בודד אך מורידים occupancy, פחות אוגרים מעלים occupancy אך מסתכנים בדליפה; אין נוסחה לנקודה המיטבית, רק מדידה.
• המטרה היא throughput מרבי, לא occupancy מרבי: חוק Little (7.4) מלמד שצריך רק מספיק warps כדי להסתיר latency, ולכן ה-kernel הכי מהיר הוא לרוב בעל occupancy בינוני - מספיק warps להסתרה, מספיק אוגרים כדי להימנע מדליפה.