לדלג לתוכן

1.4 ה Warp Scheduler ומנוע ההרצה תרגול

תרגול - ה-Warp Scheduler ומנוע ההרצה

בתרגול הזה נהפוך את עקרונות ה-warp scheduler למספרים ולניסויים מדידים. נחשב את קצב ההנפקה של ה-SM, ננסח במילים שלנו למה החלפת warp זולה פי אלף מהחלפת thread ב-CPU, נבצע אומדן חוק Little להסתרת latency, ואז נראה בעיניים - עם הcompiler ועם profiler - איך register pressure (לחץ אוגרים) מוריד את מספר ה-warps התושבים ואת ה-occupancy. עבדו לפי הסדר; חלק מהתרגילים חישוביים בלבד וחלק דורשים GPU. אפשר לעבוד על כרטיס NVIDIA מקומי או ב-Google Colab (T4 חינם). אם אתם ב-Colab, החליפו -arch=sm_90 ב--arch=sm_75 והוסיפו ! לפני כל פקודת shell. שימו לב שכמה מהמספרים (למשל 128 ליבות FP32 ל-SM, ותקרת 64 warps ל-SM) הם של Ampere/Hopper; ל-Turing (T4) יש 64 ליבות FP32 ל-SM ותקרה של 32 warps ל-SM, אז התאימו את החישובים.

הכנה - קובץ מקור ללחץ אוגרים

לתרגילים 4-7 ניצור kernel שאפשר "לנפח" את צריכת האוגרים שלו. צרו את regpressure.cu:

// regpressure.cu
#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: %s\n",                   \
                    __FILE__, __LINE__, cudaGetErrorString(err_));         \
            exit(EXIT_FAILURE);                                            \
        }                                                                  \
    } while (0)

// kernel with many live variables at once -> deliberate register pressure
__global__ void heavy(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    float x = in[i];
    float a0=x, a1=x*1.1f, a2=x*1.2f, a3=x*1.3f,
          a4=x*1.4f, a5=x*1.5f, a6=x*1.6f, a7=x*1.7f;
    #pragma unroll
    for (int k = 0; k < 32; ++k) {
        a0 = fmaf(a0, a1, a2);  a1 = fmaf(a1, a2, a3);
        a2 = fmaf(a2, a3, a4);  a3 = fmaf(a3, a4, a5);
        a4 = fmaf(a4, a5, a6);  a5 = fmaf(a5, a6, a7);
        a6 = fmaf(a6, a7, a0);  a7 = fmaf(a7, a0, a1);
    }
    out[i] = a0+a1+a2+a3+a4+a5+a6+a7;
}

int main() {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *din, *dout;
    CUDA_CHECK(cudaMalloc(&din, bytes));
    CUDA_CHECK(cudaMalloc(&dout, bytes));
    int block = 256;
    int grid  = (n + block - 1) / block;
    heavy<<<grid, block>>>(din, dout, n);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    CUDA_CHECK(cudaFree(din));
    CUDA_CHECK(cudaFree(dout));
    return 0;
}

תרגיל 1 - קצב ההנפקה של ה-SM

לפני ניסויים, נחשב את התפוקה התאורטית של מנוע ההנפקה.

  1. הריצו deviceQuery או תוכנית cudaGetDeviceProperties קצרה (מ-0.2) וקראו את מספר ה-SM (multiProcessorCount), את warpSize, ואת תדר ה-boost.
  2. בהנחה של ארבעה warp schedulers ל-SM שכל אחד מנפיק warp אחד למחזור, חשבו כמה threads מונפקים ל-SM בכל מחזור.
  3. הכפילו במספר ה-SM כדי לקבל את מספר ה-threads המונפקים בכל מחזור בכל ה-GPU. השוו את המספר הזה למספר ליבות ה-FP32 של הכרטיס (SM_count x 128 ל-Hopper). מה אתם רואים ולמה?
  4. הכפילו בתדר ה-boost (ב-Hz) כדי לקבל את קצב ההנפקה בהוראות-thread לשנייה. השוו את סדר הגודל לשיא ה-FP32 (66.9 TFLOPS ל-H100).

רמז: מספר ליבות ה-FP32 ל-SM (128 ב-Hopper) מתחלק בדיוק בארבע תת-המחיצות, ולכן כל מתזמן מזין 32 ליבות - בדיוק רוחב warp. המספרים בסעיף 3 אמורים להיות זהים; חשבו למה זו אותה כמות משני כיוונים.

תרגיל 2 - למה ההחלפה זולה פי אלף

תרגיל ניתוח, בלי קוד. נסחו במדויק את המנגנון.

  1. תארו בשתי-שלוש שורות מה בדיוק קורה כשמערכת הפעלה מחליפה thread על ליבת CPU: מה נשמר, מה נטען, ומה קורה ל-cache. אמדו את העלות במחזורים.
  2. תארו מה קורה כשה-warp scheduler מחליף warp על SM. הצביעו על שתי הסיבות שההחלפה אינה מזיזה נתונים ואינה גורמת ל-cache thrash.
  3. חשבו את היחס: אם החלפת CPU עולה ~1000 מחזורים והחלפת GPU עולה ~1 מחזור, פי כמה זולה החלפת ה-GPU? קשרו למספר ה"warps החמים" שאפשר להחזיק (עד 64 ל-SM) מול threads-על ב-CPU (1-2 לליבה).
  4. ענו: אם היינו מנסים לשמור ולשחזר את האוגרים בכל החלפת warp (כמו ב-CPU), מה זה היה עושה ליכולת של ה-GPU להחליף warp בכל מחזור?

רמז: המפתח הוא שכל ה-warps התושבים חולקים את קובץ האוגרים בו-זמנית - האוגרים שלהם כבר שם, מוקצים מראש. שמירה/שחזור היו הופכים את ההחלפה מ"שינוי מצביע" ל"העברת נתונים", והורסים את התכונה של החלפה-כל-מחזור.

תרגיל 3 - חוק Little בזעיר אנפין

נאמוד כמה warps צריך כדי להסתיר latency, ונגלה את גבולות ה-occupancy.

  1. נניח latency של טעינה מ-global memory = 400 מחזורים, ומתזמן בודד שמנפיק הוראה אחת למחזור. לפי חוק Little (מספר-בטיסה = latency x throughput), כמה הוראות עצמאיות צריכות להיות בטיסה כדי שהמתזמן לא יתבטל אף מחזור?
  2. אם כל warp תורם הוראה עצמאית אחת בכל רגע (שרשרת תלויה לגמרי), כמה warps צריך למתזמן? השוו למספר ה-warps שמתזמן בודד ב-H100 יכול להחזיק (64 warps ל-SM חלקי 4 מתזמנים).
  3. הסבירו: אם 16 warps רחוקים מ-400, איך בכל זאת kernels אמיתיים מסתירים את ה-latency? נקבו בשני המנופים.
  4. חזרו על החישוב עבור latency של פעולת חשבון = 5 מחזורים במקום 400. כמה warps צריך עכשיו למתזמן? הסבירו למה latency של חשבון קל להסתרה ו-latency של זיכרון קשה.

רמז: שני המנופים הם העלאת occupancy (עוד warps תושבים, עד התקרה) והעלאת ILP (לתת לכל warp כמה הוראות עצמאיות ברצף, למשל כמה טעינות עצמאיות לפני שצורכים תוצאה). מספר-בטיסה אפקטיבי = warps x הוראות-עצמאיות-לכל-warp.

תרגיל 4 - מדידת צריכת האוגרים

עכשיו נראה את הcompiler מדווח על האוגרים.

  1. קמפלו את regpressure.cu עם דגל הדיווח של ptxas:
    nvcc -arch=sm_90 -Xptxas -v regpressure.cu -o regpressure
  2. מצאו בפלט את השורה Used N registers עבור הפונקציה heavy. רשמו את N.
  3. חשבו ידנית, מ-N: כמה threads תושבים מקסימלית (min(2048, 65536/N)), כמה warps, ומה ה-occupancy התאורטי ב-H100.
  4. קמפלו שוב עם הגבלת אוגרים: פעם עם -maxrregcount=32, פעם עם -maxrregcount=64, ופעם עם -maxrregcount=128. לכל אחת רשמו את מספר האוגרים בפועל, את מספר ה-warps התושבים, ואת ה-occupancy. שימו לב בפלט אם הופיעו spill stores/spill loads.

רמז: -Xptxas -v מעביר את הדגל -v ל-ptxas (המרכיב ש-nvcc מפעיל כדי להפוך PTX ל-SASS). כש--maxrregcount נמוך מדי, ה-kernel לא נכנס במגבלה והcompiler "שופך" משתנים ל-local memory (שיושב ב-global) - זה ה-spill, וזה עצמו פוגע בביצועים גם אם ה-occupancy עולה.

תרגיל 5 - אימות התפוסה עם ה-API

הבה נבקש מ-CUDA עצמה לחשב לנו את התפוסה, במקום ידנית.

  1. הוסיפו ל-main (או בתוכנית נפרדת) קריאה ל-cudaOccupancyMaxActiveBlocksPerMultiprocessor, עם ה-kernel heavy, blockSize=256, ו-dynamicSMem=0.
  2. הדפיסו את מספר ה-blocks הפעילים ל-SM שהוחזר. חשבו ממנו את מספר ה-warps התושבים (blocks x 256/32) ואת ה-occupancy (חלקי 64).
  3. השוו את התוצאה לחישוב הידני מתרגיל 4 (בלי -maxrregcount). האם הם מסכימים?
  4. שנו את blockSize ל-128 ואז ל-512 והריצו שוב. האם ה-occupancy משתנה? הסבירו את הקשר בין גודל ה-block לגרנולריות של הקצאת ה-warps.

רמז: הפונקציה מקבלת מצביע ל-kernel ומחזירה כמה blocks נכנסים על SM לפי כל האילוצים (אוגרים, shared memory, warp slots) יחד. הכפלת התוצאה ב-warps-per-block ובחלוקה ב-64 נותנת את ה-occupancy התאורטי - בדיוק מה שחישבתם ביד, רק שכאן ה-driver עושה את זה.

תרגיל 6 - מקבילות מול מקביליות, חזרה ל-1.2

תרגיל קישור מושגי, בלי קוד.

  1. עבור ה-launch של regpressure (grid של 4096 blocks, block של 256), חשבו כמה threads בסך הכל בעבודה (concurrency פוטנציאלי).
  2. חשבו כמה threads תושבים בו-זמנית על SM בודד (מתרגיל 5), וכמה על כל 132 ה-SM - זהו ה-concurrency הפעיל.
  3. חשבו כמה threads רצים ממש בכל מחזור נתון (parallelism אמיתי): 128 ל-SM כפול 132.
  4. סדרו את שלושת המספרים מהגדול לקטן, והסבירו במילים שלכם: מי מהם "מקבילות" (concurrency, מנוהל לסירוגין) ומי "מקביליות" (parallelism, רץ באמת), ואיזה תפקיד ממלא ה-warp scheduler בגשר ביניהם.

רמז: העבודה הכוללת (מיליון threads) גדולה בהרבה מהתושבים בו-זמנית (מאות אלפים), שגדולים בהרבה מהרצים בפועל בכל מחזור (~16,896). ה-warp scheduler הוא מי שממיר את המאגר הגדול לתפוקה רציפה ביחידות ההרצה הצרות.

תרגיל 7 (בונוס) - פרופיילינג של הסתרת latency

נמדוד את המושגים בעולם האמיתי עם Nsight Compute.

  1. הריצו את ה-kernel תחת ncu ובקשו את מדדי ה-occupancy וה-scheduler:
    ncu --set basic ./regpressure (או ncu --metrics sm__warps_active.avg.pct_of_peak_sustained_active,smsp__issue_active.avg.pct_of_peak_sustained_active ./regpressure).
  2. אתרו את ה-Achieved Occupancy ואת ה-Issue Efficiency (אחוז המחזורים שבהם מתזמן הנפיק הוראה). רשמו אותם.
  3. קמפלו את regpressure פעם עם -maxrregcount=32 ופעם בלי, והריצו את שניהם תחת ncu. השוו את ה-Achieved Occupancy. האם העלאת ה-occupancy שיפרה את ה-Issue Efficiency, או שה-spill קילקל את הרווח?
  4. הסתכלו בקטע Warp State Statistics (או במדד smsp__average_warps_active) וזהו את סיבת ה-stall השכיחה. האם היא זיכרון או תלות חשבון? קשרו לתרגיל 3.

רמז: ncu דורש הרשאות (על מכונה מקומית ייתכן שתצטרכו sudo או הגדרת NVreg_RestrictProfilingToAdminUsers=0). ב-Colab הריצו !ncu --set basic ./regpressure. שימו לב שהעלאת occupancy אינה תמיד משפרת ביצועים - אם היא באה במחיר spills או פחות ILP, ה-Issue Efficiency יכולה דווקא לרדת. זו בדיוק הסיבה שמדידה עדיפה על ניחוש, ונחזור לזה בפרק 8.