לדלג לתוכן

5.4 ניטור וניהול NVML ו nvidia smi תרגול

תרגול - ניטור וניהול - NVML ו-nvidia-smi

בתרגול הזה תלמדו לראות את ה-GPU עובד, מבפנים. תתחילו בתשאול בסיסי עם nvidia-smi --query-gpu, אחר כך תריצו kernel שמעמיס את הכרטיס ותצפו בהספק, בטמפרטורה ובתדר משתנים בזמן אמת עם dmon. תכתבו סקריפט pynvml שדוגם את הטלמטריה תוך כדי ריצת workload ומדפיס סיכום. תשחקו עם תקרת ההספק (אם יש לכם הרשאה) ותראו במו עיניכם איך הורדת התקרה מורידה את התדר ואת התפוקה. לבסוף תסבירו למה כלים אמיתיים חייבים לפנות ל-NVML ולא לפרסר את הטקסט של nvidia-smi. עבדו לפי הסדר. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אבל הכל עובד על כל GPU של NVIDIA - רק המספרים (הספק, תדר, תקרה) יהיו של הכרטיס שלכם.

הכנה

התקינו את binding ה-Python של NVML ואמתו שהוא עובד:

pip install nvidia-ml-py
python -c "import pynvml; pynvml.nvmlInit(); print('NVML OK'); pynvml.nvmlShutdown()"

צרו kernel חישובי כבד שיעמיס את הכרטיס מספיק זמן כדי שנספיק לצפות בו. burn.cu:

#include <cstdio>
#include <cstdlib>

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

// compute-bound kernel: dense FMA loop, no memory accesses in the loop
__global__ void burn(float* out, int iters) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float a = idx * 0.5f, b = 1.0001f, c = 0.9999f;
    for (int i = 0; i < iters; i++) {
        a = a * b + c;      // FMA
        a = a * c + b;      // FMA
    }
    out[idx] = a;           // single write at the end, so the compiler doesn't eliminate the loop
}

int main(int argc, char** argv) {
    int secs = (argc > 1) ? atoi(argv[1]) : 15;     // roughly how many seconds to load
    const int blocks = 4096, threads = 256;
    const int n = blocks * threads;
    float* d_out;
    CUDA_CHECK(cudaMalloc(&d_out, (size_t)n * sizeof(float)));

    // run repeatedly to keep the card busy for ~secs seconds
    for (int rep = 0; rep < secs * 20; rep++) {
        burn<<<blocks, threads>>>(d_out, 2'000'000);
    }
    CUDA_CHECK(cudaDeviceSynchronize());
    CUDA_CHECK(cudaGetLastError());
    printf("done\n");
    cudaFree(d_out);
    return 0;
}

הcompilation והרצה:

nvcc -arch=sm_90 burn.cu -o burn      # replace sm_90 with your architecture
./burn 15                             # loads the card for roughly 15 seconds

תרגיל 1 - תשאול בסיסי עם query-gpu

  1. הריצו nvidia-smi יבש והכירו את הטבלה. זהו את השדות: Temp, Perf, Pwr:Usage/Cap, Memory-Usage, GPU-Util, Persistence-M, Compute M..
  2. תשאלו בדיוק את השדות שאתם רוצים בפורמט CSV:
    nvidia-smi --query-gpu=name,power.draw,power.limit,temperature.gpu,utilization.gpu,memory.used,memory.total,clocks.sm,clocks.mem,pstate --format=csv
  3. הריצו שוב עם --format=csv,noheader,nounits וראו את ההבדל בפלט. איזו גרסה נוחה יותר לפרסור בסקריפט?
  4. גלו את הטווח המותר של תקרת ההספק לכרטיס שלכם: nvidia-smi -q -d POWER וחפשו Min Power Limit ו-Max Power Limit. מה תקרת ברירת המחדל?

רמז: --query-gpu מקבל רשימת שדות מופרדת בפסיקים ו---format הכרחי. השדות הנפוצים: power.draw, power.limit, temperature.gpu, utilization.gpu, memory.used, clocks.sm, clocks.mem, pstate. ל-H100 ברירת המחדל של התקרה היא 700W והטווח בערך 200-700W; לכרטיס שלכם זה יהיה שונה. noheader,nounits נותן ערכים נקיים בלי כותרת ובלי היחידות - הכי נוח ל-awk/פרסור.

תרגיל 2 - צפייה חיה בהספק ותדר עם dmon

  1. פתחו שני טרמינלים. באחד הריצו את ה-burn לזמן ארוך: ./burn 30.
  2. בשני, מיד, הריצו טלמטריה חיה:
    nvidia-smi dmon -s puct -d 1
  3. צפו בעמודות pwr (הספק), gtemp (טמפרטורה), sm (ניצול), pclk (תדר SM) לאורך זמן. תעדו את השורות: מה קורה בשנייה הראשונה מול השנייה העשירית?
  4. במקביל, בטרמינל שלישי, הריצו nvidia-smi pmon -s um -d 1 וזהו את התהליך burn ואת ה-PID שלו.
  5. בזמן שה-burn רץ בשיא, הריצו nvidia-smi -q -d PERFORMANCE. האם SW Power Cap הוא Active? מה זה אומר על ה-workload?

רמז: dmon -s puct בוחר power/temp (p), utilization (u), clocks (c), PCIe (t); -d 1 = דגימה כל שנייה. שימו לב לתבנית: בהתחלה התדר קופץ לבוסט המלא, ואז - ככל שההספק נוגע בתקרה והטמפרטורה עולה - התדר יורד. אם SW Power Cap : Active, הכרטיס מווסת בגלל תקרת הספק, כלומר ה-workload power-bound. pmon מראה per-process; חפשו את השם burn בעמודת command.

תרגיל 3 - סקריפט pynvml שדוגם workload

  1. כתבו sampler.py שמדגם הספק וטמפרטורה בלולאה כל 100ms, במשך מספר שניות, ובסוף מדפיס מינימום, מקסימום וממוצע לכל אחד. השתמשו ב-nvmlDeviceGetPowerUsage (מגיע ב-mW), nvmlDeviceGetTemperature, nvmlDeviceGetUtilizationRates, nvmlDeviceGetClockInfo.
  2. הריצו את ה-burn בטרמינל אחד (./burn 20), ובמקביל הריצו את sampler.py בשני.
  3. הדפיסו סיכום בסגנון: power avg/max, temp avg/max, sm util avg, sm clock min/max. מה הפער בין תדר ה-SM המקסימלי (בהתחלה) למינימלי (אחרי חימום)?
  4. (הרחבה) הוסיפו קריאה ל-nvmlDeviceGetPowerManagementLimit והדפיסו את יחס ה-avg_power / limit. מה הוא מלמד אתכם על עד כמה ה-workload דוחף בתקרה?

רמז: הסקלד: nvmlInit() -> nvmlDeviceGetHandleByIndex(0) -> לולאה עם time.sleep(0.1) שצוברת ערכים לרשימות -> nvmlShutdown() ב-finally. זכרו את היחידות: הספק מגיע ב-מיליוואט (חלקו ב-1000 ל-W), הזיכרון ב-בתים. nvmlDeviceGetUtilizationRates מחזיר אובייקט עם .gpu ו-.memory. nvmlDeviceGetClockInfo(h, pynvml.NVML_CLOCK_SM) נותן את תדר ה-SM ב-MHz. עטפו את הכל ב-try/finally ותמיד קראו nvmlShutdown().

תרגיל 4 - תקרת הספק והשפעתה על התפוקה

  1. תשאלו את התקרה הנוכחית ואת הטווח המותר:
    nvidia-smi -q -d POWER | grep -i "power limit"
  2. הריצו את ה-burn פעם אחת עם התקרה המלאה ומדדו כמה זמן הוא לוקח (time ./burn 5 - או הוסיפו תזמון פנימי ב-CUDA events). תעדו גם את התדר הממוצע מ-dmon.
  3. אם יש לכם הרשאת root, הורידו את התקרה למחצית בערך והריצו שוב:
    sudo nvidia-smi -pm 1              # persistence mode
sudo nvidia-smi -pl 350           # lower limit (adjust to your range)
time ./burn 5
nvidia-smi dmon -s pc -c 5        # 5 samples of power and clock
  4. השוו: מה קרה לתדר ה-SM? מה קרה לזמן הריצה של ה-burn? הסבירו את הקשר בין תקרת הספק, תדר ותפוקה.
  5. החזירו את התקרה לברירת המחדל: sudo nvidia-smi -pl 700 (או הערך המקורי שלכם).

רמז: אם אין לכם הרשאת root (נפוץ בענן משותף), הפקודה תיכשל ב-Insufficient Permissions - זה תקין, פשוט תארו מה הייתם מצפים לראות: תקרה נמוכה יותר -> הכרטיס נוגע בה מוקדם יותר -> SW Power Cap : Active -> התדר יורד -> אותו kernel לוקח יותר זמן. תקרת ההספק היא תקציב: פחות ואטים = פחות מחזורי שעון לשנייה = פחות FLOP לשנייה. זכרו להחזיר את התקרה בסוף.

תרגיל 5 - למה NVML ולא טקסט של nvidia-smi

  1. כתבו סקריפט "רע" שמפרסר את הפלט הטקסטואלי של nvidia-smi כדי לחלץ את ההספק - למשל nvidia-smi | grep W | awk .... גרמו לו לעבוד.
  2. כתבו סקריפט "טוב" שמחלץ את אותו ערך דרך pynvml (nvmlDeviceGetPowerUsage).
  3. הריצו את שניהם ואמתו שהם מסכימים.
  4. הסבירו בכתב: מדוע הסקריפט הראשון שביר? מה יכול לשבור אותו בלי שום שינוי בקוד שלכם? מהי החלופה המובנית הביניים (בין טקסט חופשי ל-NVML) ש-nvidia-smi מציע?

רמז: הסקריפט הטקסטואלי מסתמך על מיקום מדויק, רוחב עמודות ומחרוזות בפלט של nvidia-smi - וכל אלה אינם מובטחים כיציבים בין גרסאות driver. עדכון driver יכול לשנות סדר שדות או פורמט וישבור את הפרסור, בלי שנגעתם בקוד. החלופה המובנית הביניים היא nvidia-smi -x (XML) או --query-gpu ... --format=csv; החלופה הטובה ביותר היא ה-C API של NVML ישירות (או pynvml/nvml_wrapper) - זה ה-API שהמקור מבטיח את יציבותו.

תרגיל 6 (בונוס) - exporter קטן וזיהוי throttling

  1. הרחיבו את sampler.py כך שידגום כל שנייה ברצף (לא לזמן קצוב), ובכל דגימה יבדוק את סיבות הוויסות דרך nvmlDeviceGetCurrentClocksThrottleReasons(h) וישווה מול הקבועים nvmlClocksThrottleReasonSwPowerCap ו-nvmlClocksThrottleReasonHwThermalSlowdown.
  2. הדפיסו שורה לכל שנייה בפורמט CSV: timestamp,power_w,temp_c,sm_clock,throttle. כשה-burn רץ, האם אתם רואים את SwPowerCap נדלק?
  3. (מחשבה) קשרו לפרק 7.1: אם ה-workload power-bound (התדר יורד כשההספק נוגע בתקרה), מה זה אומר על עד כמה אופטימיזציה של הקוד תעזור? מה כן יעזור?
  4. (מחשבה) איך היה נראה exporter כזה בפרודקשן - למשל שמזין Prometheus? מדוע דווקא ל-NVML יש כאן יתרון על פני nvidia-smi?

רמז: nvmlDeviceGetCurrentClocksThrottleReasons מחזיר מסכת סיביות; בדקו כל סיבה עם reasons & pynvml.nvmlClocksThrottleReasonSwPowerCap. אם ה-bit של SwPowerCap דולק תחת עומס, ה-workload power-bound - ואז אופטימיזציית קוד לא תשבור את תקרת ה-700W; רק העלאת התקציב (הספק/קירור) או הפחתת עבודה תעזור. זה בדיוק הקשר לפרק 7.1: תפוקת השיא בפועל נקבעת על ידי power ingress ו-heat egress, לא רק על ידי יעילות הקוד. ל-exporter בפרודקשן NVML עדיף כי ה-API שלו יציב וקריאותיו זולות, בעוד פרסור טקסט של nvidia-smi שביר ומריץ תת-תהליך בכל דגימה.