לדלג לתוכן

7.4 Overhead, latency hiding וחוק Little תרגול

תרגול - Overhead, latency hiding וחוק Little

בתרגול הזה תמדדו במו ידיכם את שלושת הרעיונות של השיעור. תתחילו במדידת overhead הlaunch של kernel ריק לגמרי - כדי לבודד את זמן הhost מכל חישוב - ותקבלו את המספר של ~10 מיקרו-שניות לlaunch בעצמכם. אחר כך תיישמו את חוק Little בעיפרון: תחשבו כמה פעולות צריך "בטיסה" וכמה warps צריך כדי להסתיר latency נתון, ותתרגמו את זה ל-occupancy על H100. משם תמירו את לולאת הlaunches לגרף CUDA ותכמתו כמה overhead הוסרה, תחשבו את ה-occupancy הדרוש לתרחיש קונקרטי, תסבירו את תוצאת Volkov המפורסמת (30 מול 24 warps), ולבסוף תראו את הסתרת ה-latency קורית בפועל במיקרו-בנצ'מרק. עבדו לפי הסדר - כל תרגיל בונה על קודמו. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אבל כל תרגיל כאן ירוץ על כל GPU של NVIDIA (רק המספרים המוחלטים ישתנו).

הכנה

צרו קובץ overhead.cu. לאורך התרגול נשתמש ב-kernel ריק לגמרי - noop - כדי שזמן הdevice יהיה אפסי וכל מה שנמדד יהיה host overhead טהורה. הקומפילציה (החליפו את sm_90a בארכיטקטורה שלכם - למשל sm_75 ל-T4):

nvcc -O2 -arch=sm_90a -o overhead overhead.cu

מקרו הבדיקה וה-kernel הריק:

#include <cstdio>
#include <cstdlib>
#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)

__global__ void noop() { }   // zero work - all the time we measure is host overhead

תרגיל 1 - מדידת overhead הlaunch של kernel ריק

  1. צרו stream לא-default עם cudaStreamCreate.
  2. כתבו לולאה שמשגרת noop<<<1, 1, 0, stream>>>() בדיוק N = 100000 פעמים, ואחריה cudaStreamSynchronize. עטפו את כל הלולאה במדידת cudaEvent (record לפני, record אחרי, cudaEventElapsedTime).
  3. הריצו איטרציית חימום אחת של כל הלולאה לפני המדידה (כדי לספוג את ה-JIT החד-פעמי ואת אתחול ההקשר).
  4. חשבו ודווחו את הoverhead לכל launch: זמן_כולל_במיקרו / N.
  5. חזרו על המדידה גם עם noop<<<1, 1>>>() על ה-default stream (בלי הפרמטר stream). האם המספר משתנה מהותית?

רמז: מכיוון שה-kernel ריק, זמן הdevice שלו כמעט אפס, ולכן כמעט כל הזמן שתמדדו הוא הhost overhead - אריזת הארגומנטים (כאן אין ארגומנטים), תצורת הlaunch, והפנייה לדרייבר. אתם אמורים לקבל מספר בסדר גודל של יחידות מיקרו-שניות (בערך 3-10 us), בדיוק המספר מההרצאה. השתמשו ב-cudaEvent (זמן GPU) ולא בשעון CPU. הריצו הרבה launches (מאה אלף) כדי שהחלוקה תהיה יציבה.

תרגיל 2 - חוק Little בעיפרון

בלי לכתוב קוד, חשבו ורשמו את התשובות (השתמשו במחשבון). זהו התרגיל התיאורטי המרכזי.

  1. פעולות בטיסה. נתון latency של גישה לזיכרון = 400 מחזורים, ו-throughput רצוי = פקודה אחת למחזור. כמה פעולות חייבות להיות בטיסה בו-זמנית לפי חוק Little?
  2. throughput גבוה יותר. אותו latency (400 מחזורים), אבל עכשיו throughput רצוי = 10 פקודות למחזור. כמה פעולות בטיסה עכשיו? מה היחס לתשובה מ-1, ולמה?
  3. מ-פעולות ל-warps. ה-warp מנפיק פקודה עבור 32 threads יחד. המירו את תשובת סעיף 1 (400 פעולות) למספר warps.
  4. שרשרת תלויה. נתונה שרשרת ה-SASS מההרצאה: LDG (400) -> IMUL (6) -> IADD (4) -> IMUL (6), סה"כ latency. חשבו את ה-latency הכולל, ואז את מספר ה-warps הדרוש להסתרתו המלאה (חלקו ב-32).
  5. תרגום ל-occupancy. על H100 המקסימום הוא 64 warps ל-SM. מה ה-occupancy (באחוזים) שמתאים למספר ה-warps מסעיף 4? האם צריך occupancy גבוה כדי להסתיר את השרשרת הזו?

רמז: הנוסחה היא concurrency = latency x throughput. לסעיף 3 ו-4 חלקו את מספר הפעולות/המחזורים ברוחב ה-warp (32). לסעיף 5 חלקו את מספר ה-warps ב-64 והכפילו ב-100. שימו לב לתובנה: המספר שתקבלו קטן יחסית - כלומר occupancy צנוע כבר מסתיר שרשרת load אופיינית, וזה בדיוק למה "יותר occupancy" הוא לא תמיד המטרה (נעמיק בפרק 8).

תרגיל 3 - המרת הלולאה לגרף וכימות הoverhead שהוסרה

  1. קחו את לולאת ה-N הlaunches מתרגיל 1, ולכדו אותה לגרף CUDA (כפי שראינו בשיעור 5.5): cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal), הלולאה של N launches noop על ה-stream, ואז cudaStreamEndCapture(stream, &graph).
  2. הגישו למכונה עם cudaGraphInstantiate, ומדדו עם cudaEvent את הזמן של cudaGraphLaunch בודד (עם cudaStreamSynchronize אחריו), אחרי חימום, מינימום על 50 חזרות.
  3. השתמשו ב-N = 512 (לא 100000 - הגרף מחזיק צומת לכל launch, ו-512 מספיק להדגמה). מדדו גם את הלולאה הרגילה עם אותו N.
  4. דווחו שלושה מספרים: זמן הלולאה, זמן ה-replay, וה-speedup. כמה מיקרו-שניות של host overhead הוסרו לכל הרצה?
  5. הסבירו: הגרף לא מאיץ את ה-kernels (הם ריקים ממילא) - אז מה בדיוק נחסך?

רמז: בתוך הלכידה אסור להשתמש בשום פעולה חוסמת. הגרף ממיר את N הlaunches הנפרדות לקריאת host אחת, ולכן זמן ה-replay כמעט אינו תלוי ב-N. הoverhead שהוסרה היא בערך (זמן_הלולאה - זמן_ה-replay). מכיוון שה-kernels ריקים, כאן זה הניסוי הנקי ביותר לבידוד הhost overhead לבדה.

תרגיל 4 - חישוב ה-occupancy הדרוש לתרחיש נתון

בלי קוד. נתון kernel ממשי שבו שרשרת התלות הארוכה ביותר של כל thread היא: שני loads מהזיכרון הגלובלי (כל אחד 400 מחזורים) שאינם תלויים זה בזה, ואחריהם רצף אריתמטי תלוי באורך 40 מחזורים.

  1. מהו ה-latency שחייבים להסתיר? שימו לב: שני ה-loads בלתי-תלויים - warp בודד יכול להנפיק את שניהם לפני שהוא נתקע. כמה פעולות בטיסה warp בודד מספק כאן (ILP)?
  2. חשבו את ה-latency האפקטיבי שנשאר להסתרה על-ידי warps נוספים, בהתחשב ב-ILP של סעיף 1.
  3. חשבו את מספר ה-warps הדרוש, ותרגמו ל-occupancy על H100 (מתוך 64).
  4. השוו לתרחיש חלופי שבו שני ה-loads כן תלויים זה בזה (שרשרת אחת ארוכה). כמה warps צריך עכשיו? מה תפקיד ה-ILP במסקנה?

רמז: ILP (מקביליות ברמת הפקודה) ו-occupancy הם שני מנופים חלופיים לאותו יעד - להחזיק מספיק פעולות בטיסה. כשיש loads בלתי-תלויים, warp בודד מחזיק כמה בטיסה, וצריך פחות warps. כשהכל שרשרת תלויה אחת, כל warp מחזיק פקודה אחת בטיסה, וצריך יותר warps. חשבו כמה פעולות בטיסה כל warp תורם, וכמה warps צריך כדי להגיע ל-concurrency הכולל שחוק Little דורש.

תרגיל 5 - הסבר תוצאת Volkov (30 מול 24)

בלי קוד. תוצאת Volkov: צריך בערך 30 warps להסתיר latency של זיכרון, ובערך 24 warps להסתיר latency אריתמטי - למרות ש-latency הזיכרון גדול בהרבה.

  1. נסחו במילים למה זה מפתיע: מה הייתה הציפייה הנאיבית, בהינתן שה-latency של זיכרון הוא מאות מחזורים וה-latency האריתמטי הוא יחידות מחזורים?
  2. כתבו את חוק Little עבור שני המקרים בנפרד, וסמנו לכל אחד אם ה-latency גבוה/נמוך ואם ה-throughput גבוה/נמוך.
  3. הסבירו מדוע המכפלה (ה-concurrency) יוצאת דומה בשני המקרים.
  4. קשרו לרעיון "קיר הזיכרון" מ-7.1 וליחס Tensor Core:CUDA Core של בערך 100:1: איזה גורם בחוק Little משקף את הפער הזה?
  5. הסיקו מדוע זו תכונה שימושית של המכונה עבור kernel שמערבב אריתמטיקה וזיכרון.

רמז: concurrency = latency x throughput. לזיכרון: latency גבוה, throughput נמוך (רוחב-פס HBM נמוך). לאריתמטיקה: latency נמוך, throughput גבוה (רוחב-פס אריתמטי גבוה, פי ~100 עם Tensor Cores). המכפלות מתקזזות. המסקנה השימושית: אותה תצורת occupancy מסתירה עומס מעורב.

תרגיל 6 (בונוס) - לראות את הסתרת ה-latency קורית

  1. כתבו kernel שכל thread שלו מבצע שרשרת תלויה ארוכה של גישות זיכרון (למשל pointer-chasing או רצף a = data[a] שחוזר הרבה פעמים), כך שה-latency של הזיכרון שולט.
  2. השיקו אותו עם מספר threads-per-block עולה: 32, 64, 128, 256, 512, 1024 (כלומר 1, 2, 4, 8, 16, 32 warps ל-block), על block בודד או על מעט blocks, ומדדו את הזמן הכולל או את ה-throughput האפקטיבי (גישות לשנייה).
  3. שרטטו (או תארו) את ה-throughput כפונקציה של מספר ה-warps. היכן הוא מתייצב (מגיע לרמה)?
  4. קשרו את נקודת ההתייצבות לחוק Little: מספר ה-warps שבו ה-throughput מפסיק להשתפר הוא בערך המספר הדרוש להסתרת ה-latency. השוו לחישוב התיאורטי מתרגיל 2.
  5. (רשות) הריצו ncu --set full ./overhead על ה-kernel והסתכלו במדד "Achieved Occupancy" ובסיבות ה-stall (למשל Stall Long Scoreboard) - הם מספרים את אותו סיפור.

רמז: כשמעט warps רצים, ה-warp scheduler נשאר בלי warp כשיר בזמן שכולם תקועים על loads - ה-throughput נמוך. ככל שמוסיפים warps, יש יותר מה להנפיק בזמן שאחרים תקועים, וה-throughput עולה - עד שהוא רווי: הגעתם למספר ה-warps שמסתיר את כל ה-latency, ומעבר לזה אין שיפור. זו הדגמה חיה של חוק Little ושל הנקודה "מספיק occupancy, לא מקסימלי". השתמשו ב-cudaEvent למדידה, עם חימום ומינימום על כמה חזרות.