לדלג לתוכן

9.1 פרופיילינג Nsight Systems ו Nsight Compute תרגול

תרגול - Nsight Systems ו-Nsight Compute

בתרגול הזה תפעילו את שני ה-profilers על קוד אמיתי ותקראו את הפלט. תבנו אפליקציה קטנה overhead-bound, תפרפלו אותה עם nsys ותמצאו את הפער ואת חוסר החפיפה בציר הזמן; תריצו kernel אחד תחת ncu עם --set full ותקראו את קטעי ה-roofline, ה-occupancy, הזיכרון וה-stalls; תאשרו עם ncu בעיית coalescing ובעיית bank conflict מהשיעורים ב-8.4; ולבסוף תקשרו שורה חמה ב-SASS למקור עם -lineinfo. עבדו לפי הסדר - כל תרגיל בונה על הקודם. אפשר לעבוד על כל כרטיס NVIDIA מקומי או בחינם ב-Google Colab. אם אתם ב-Colab על T4, החליפו -arch=sm_90 ב--arch=sm_75, הוסיפו ! לפני כל פקודת shell, וזכרו שנקודת הרכס ורוחבי הפס יהיו של ה-T4 ולא של ה-H100 - העיקרון והדוחות זהים, המספרים המוחלטים שונים.

הכנה - כלים ומאקרו

לכל התרגילים נשתמש במאקרו בדיקת השגיאות הסטנדרטי; שימו אותו בראש כל קובץ .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 %s at %s:%d\n",                    \
                    cudaGetErrorString(err_), __FILE__, __LINE__);         \
            exit(EXIT_FAILURE);                                            \
        }                                                                  \
    } while (0)

ודאו שהכלים מותקנים: nsys --version ו-ncu --version (שניהם מגיעים עם ה-CUDA Toolkit). על שרת ללא GUI נעבוד עם פלט טקסטואלי בלבד - nsys stats ו-ncu לטרמינל - בלי לפתוח את ה-GUI. הערה על הרשאות: ncu דורש גישה למוני החומרה; אם אתם מקבלים ERR_NVGPUCTRPERM, הריצו עם sudo או בקשו מהמנהל להפעיל את הדגל NVreg_RestrictProfilingToAdminUsers=0.

תרגיל 1 - nsys על אפליקציה overhead-bound

נבנה אפליקציה שמשגרת המון kernels זעירים, ונראה את הפערים בציר הזמן.

  1. צרו overhead.cu עם kernel זעיר (למשל tiny_add שמוסיף קבוע לאלמנט אחד לכל thread) על מערך קטן מאוד (למשל N = 1024).
  2. ב-main, שגרו את ה-kernel בלולאה 10,000 פעם, בזה אחר זה, עם cudaDeviceSynchronize() בסוף.
  3. קמפלו nvcc -arch=sm_90 -O3 overhead.cu -o overhead והריצו nsys profile -o rep_overhead ./overhead.
  4. הריצו nsys stats rep_overhead.nsys-rep והסתכלו בטבלאות: כמה זמן הלך לקריאות ה-CUDA API (cudaLaunchKernel) לעומת זמן ה-GPU בפועל? מה היחס?
  5. הסבירו: האם האפליקציה overhead-bound, compute-bound או memory-bound? מה המספר ~10 מיקרו-שנייה לdispatch (שיעור 7.4) אומר כאן?

רמז: באפליקציה כזו, זמן ה-CPU בקריאות cudaLaunchKernel יעלה בהרבה על זמן ה-GPU - זהו התסמין הקלאסי של overhead-bound. חפשו בטבלת "CUDA API Summary" את cudaLaunchKernel ובטבלת "CUDA GPU Kernel Summary" את זמן ה-kernel. 10,000 dispatches כפול ~10 מיקרו-שנייה = ~100 מילישנייה של overhead בלבד.

תרגיל 2 - מציאת חוסר חפיפה בין copy ל-compute

נבנה pipeline של העתקה-וחישוב ונראה מתי הם לא חופפים.

  1. צרו pipeline.cu שמעבד M אצוות (chunks) של נתונים: לכל אצווה, cudaMemcpy HtoD, ואז kernel, ואז cudaMemcpy DtoH - הכל על ה-default stream, בטור.
  2. פרפלו עם nsys profile -o rep_serial ./pipeline וקראו את ציר הזמן (או nsys stats). ודאו שההעתקות וה-kernels רצים בזה-אחר-זה, בלי חפיפה.
  3. שכתבו לגרסה עם streams מרובים ו-pinned memory (cudaHostAlloc), כך שההעתקה של אצווה אחת תחפוף לחישוב של אצווה אחרת (שיעור 5.3). השתמשו ב-cudaMemcpyAsync.
  4. פרפלו את הגרסה החדשה (rep_overlap) והשוו: האם עכשיו רואים HtoD של אצווה חופף ל-kernel של אצווה קודמת? מה קרה לזמן הכולל?

רמז: cudaMemcpy רגיל (סינכרוני) על זיכרון pageable לא יחפוף לכלום. צריך שלושה דברים יחד: cudaHostAlloc (pinned), cudaMemcpyAsync, ו-streams שונים לאצוות שונות. ב-nsys stats חפשו בשורות ה-Memory Operations וה-Kernels את חותמות הזמן - בגרסה הטובה טווחי הזמן שלהם יחפפו.

תרגיל 3 - ncu --set full וקריאת הקטעים

ניקח kernel אחד ונצלול אליו לעומק עם ncu.

  1. השתמשו ב-rowsum.cu משיעור 8.4 (עם rowsum_bad ו-rowsum_good), או כתבו kernel פשוט של SAXPY (y = a*x + y) על מערך גדול (N = 64M).
  2. קמפלו עם -lineinfo: nvcc -arch=sm_90 -O3 -lineinfo saxpy.cu -o saxpy.
  3. הריצו ncu --set full -k saxpy --launch-count 1 -o rep_saxpy ./saxpy.
  4. קראו את הקטעים בפלט (או ב-ncu -i rep_saxpy.ncu-rep --page details): מה קטע ה-roofline / Speed of Light אומר - compute-bound או memory-bound? מה ה-arithmetic intensity המדוד, וכיצד הוא מתייחס לנקודת הרכס 295 של ה-H100 (שיעור 7.2)?
  5. הסתכלו בקטע ה-Occupancy: מה ה-occupancy התיאורטי מול המושג? בקטע ה-Warp State / Stalls: מהי סיבת העצירה הדומיננטית?

רמז: SAXPY הוא הדוגמה הקלאסית לmemory-bound - arithmetic intensity של 1/6 FLOP/byte, הרבה מתחת ל-295, אז הנקודה תשב על הגג האלכסוני. סיבת העצירה הדומיננטית תהיה long scoreboard (המתנה ל-LDG מ-global). ה-occupancy יהיה גבוה יחסית כי SAXPY משתמש במעט registers.

תרגיל 4 - אישור בעיית coalescing עם ncu

נחזור ל-kernel מ-8.4 ונאשר את ה-coalescing ישירות עם המדד.

  1. קמפלו את rowsum.cu (rowsum_bad מול rowsum_good משיעור 8.4).
  2. הריצו כל אחד תחת ncu עם המדד: ncu --metrics l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio ./rowsum.
  3. איזו גרסה מראה ~4 sectors לבקשה ואיזו ~32? קשרו למספר: 128B transaction = 4 sectors של 32B (מאוחד), עד 32 (מפוזר לגמרי).
  4. הריצו את שתי הגרסאות תחת ncu --set full והסתכלו בקטע ה-Memory Workload. כיצד ה-sectors-per-request מסביר את ה-roofline של כל גרסה (מרחק מהגג האלכסוני)?
  5. הצליבו: האם קטע ה-Stalls של הגרסה הרעה מראה יותר long scoreboard stalls מהגרסה הטובה? למה זה עקבי עם ה-sectors?

רמז: rowsum_bad נותן ל-threads שכנים כתובות במרחק N*4 בייט (32 sectors לבקשה); rowsum_good נותן להם כתובות רצופות (4 sectors). יותר sectors מיותרים = יותר תעבורת global = יותר LDG = יותר long scoreboard stalls. שלושת המדדים (sectors, roofline, stalls) מספרים את אותו סיפור מזוויות שונות.

תרגיל 5 - אישור bank conflict עם ncu

נאשר את ההתנגשות ה-32-way של ה-transpose משיעור 8.4.

  1. קמפלו את transpose.cu משיעור 8.4 עם transposeNaive (tile[32][32]) ו-transposePadded (tile[32][33]).
  2. הריצו כל גרסה תחת ncu: ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum ./transpose.
  3. כמה התנגשויות מראה הגרסה הנאיבית וכמה המרופדת? הראו שהמרופדת צנחה כמעט לאפס.
  4. בקטע ה-Memory Workload של --set full, איפה מופיעה הבעיה - במפלס ה-global או ה-shared? הסבירו למה ה-global דווקא נקי בשתי הגרסאות (רמז: שיעור 8.4, הקריאה/כתיבה ל-global מאוחדות).
  5. בקטע ה-Stalls של הגרסה הנאיבית - האם רואים short scoreboard גבוה (המתנה ל-LDS מ-shared)? קשרו לסריאליזציה של ה-bank conflict.

רמז: transposeNaive מייצר מיליוני התנגשויות; transposePadded מייצר ~0. ה-bank conflict הוא בקריאה מ-shared (LDS), ולכן הוא short scoreboard, לא long - זה מבדיל אותו מבעיית ה-coalescing מתרגיל 4 שהיא long scoreboard על global.

תרגיל 6 - קורלציית מקור-SASS

נאתר את השורה החמה בקוד דרך תצוגת ה-Source של ncu.

  1. ודאו שקמפלתם עם -lineinfo (בלי זה אין מיפוי מקור). הריצו ncu --set full --import-source yes -k rowsum_bad -o rep_src ./rowsum.
  2. פתחו את תצוגת ה-Source (ב-GUI, או ncu -i rep_src.ncu-rep --page source). מיינו את השורות לפי מספר ה-stalls / ה-sampling.
  3. איזו שורת מקור היא ה"חמה" ביותר (הכי הרבה מחזורי stall)? איזו הוראת SASS מקורה?
  4. הסבירו: מדוע השורה acc += A[row*N+col] היא החמה, ומדוע הוראת ה-LDG.E שלה עוצרת על long scoreboard? קשרו לכל השרשרת: מקור -> SASS -> stall reason -> roofline.

רמז: השורה החמה תהיה הגישה ל-global (A[row*N+col]), וה-SASS שלה LDG.E.SYS. הוראה זו כותבת scoreboard וממתינה לזיכרון שחוזר מ-HBM, ולכן היא long scoreboard stall - בדיוק מה שראינו ב-8.3. אם השורה החמה היא הגישה מ-shared, ה-SASS יהיה LDS וה-stall יהיה short.

תרגיל 7 (בונוס) - לולאת האופטימיזציה המלאה

תרגיל אינטגרציה שמפעיל את כל השרשרת.

  1. קחו את pipeline.cu הגרוע מתרגיל 2 (default stream, בלי חפיפה, עם kernel כבד יחסית).
  2. צעד nsys: פרפלו וקבעו - האם הבעיה הראשונה היא overhead/חוסר overlap (בעיה מערכתית) או שה-kernel עצמו איטי? תקנו את המערכתי קודם (streams + pinned) ומדדו שוב.
  3. צעד ncu: אחרי שסגרתם את החפיפה, פרפלו את ה-kernel עצמו עם --set full. הcompute-bound או memory-bound? מה הצוואר - coalescing? occupancy? bank conflict?
  4. תקנו את הצוואר של ה-kernel, וחזרו ל-nsys. מה הצוואר החדש? כתבו טבלה: בכל איטרציה - איזה כלי השתמשתם, מה מצאתם, מה תיקנתם, ומה ה-speedup.
  5. הסבירו במשפט למה סדר הכלים חשוב: מה היה קורה אם הייתם מתחילים מ-ncu לפני nsys?

רמז: זו מתודולוגיית ה-Theory of Constraints משיעור 7.1 - מצא צוואר, הרם, חזור על החדש. אם מתחילים מ-ncu, מסכנים לאפטמז kernel שבכלל לא צוואר הבקבוק - כל הזמן אולי הולך לoverhead או להעתקות שלא חופפות, ו-ncu (שרואה kernel בודד) לעולם לא יראה את זה. תמיד nsys ("איפה") לפני ncu ("למה").