לדלג לתוכן

1.6 היררכיית הזיכרון בחומרה תרגול

תרגול - היררכיית הזיכרון בחומרה

בתרגול הזה תמדדו את פירמידת הזיכרון של ה-GPU במקום רק לצייר אותה. תשלפו את המספרים האמיתיים של הכרטיס שלכם, תמדדו במו ידיכם עד כמה גישה ל-global memory איטית מפעולת חישוב, תראו איפה אוגר שגלש חי באמת ומה הוא עולה, ותסבירו במדויק למה cache ה-GPU שונה במהותו מcache CPU. עבדו לפי הסדר - כל תרגיל בונה על הקודם. אפשר לעשות את הכל על כרטיס NVIDIA מקומי או בחינם ב-Google Colab. אם אתם ב-Colab, החליפו -arch=sm_90 ב--arch=sm_75 (ה-compute capability של ה-T4), הוסיפו ! לפני כל פקודת shell, וזכרו שלכל כרטיס יש מספרים אחרים - העיקרון הוא שנשמר.

הכנה - קובץ המדידה

לאורך התרגול נשתמש בקובץ מדידה אחד. צרו אותו עכשיו; נשלים ונשנה אותו תרגיל אחרי תרגיל.

// memlat.cu
#include <cstdio>
#include <cstdlib>
#include <cstdint>
#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)

int main() {
    // filled in during exercises 2 and 3
    return 0;
}

קמפלו תמיד עם ציון הארכיטקטורה, למשל nvcc -arch=sm_90 memlat.cu -o memlat.

תרגיל 1 - ציור הפירמידה מהמספרים של הכרטיס שלכם

לפני שמודדים, שולפים את הממדים.

  1. כתבו תוכנית קצרה (או הרחיבו את memlat.cu) שקוראת ל-cudaGetDeviceProperties ומדפיסה: multiProcessorCount, regsPerMultiprocessor, sharedMemPerMultiprocessor, l2CacheSize, ו-totalGlobalMem.
  2. חשבו ידנית: מה גודל קובץ האוגרים לכל SM ב-KiB (רמז: כל אוגר 4 בייט)? מה הגודל המצטבר של קובץ האוגרים ושל ה-L1/shared על פני כל ה-SM?
  3. ציירו את הפירמידה בארבע מדרגות (אוגרים, L1/shared, L2, GPU RAM). לכל מדרגה רשמו: גודל ל-SM (או device-wide), הגודל המצטבר, הטכנולוגיה (SRAM/DRAM), והמנהל (מתכנת/חומרה).
  4. לצד כל מדרגה כתבו את המהירות היחסית בסדרי גודל: אוגרים = בסיס, L1 = פי 10 איטי, וכן הלאה. אין צורך במספרים מדויקים - רק בסדר גודל.

רמז: regsPerMultiprocessor מוחזר כמספר אוגרים (לא בייטים) - הכפילו ב-4 ל-בייטים. l2CacheSize מוחזר בבייטים; חלקו ב-1024*1024 ל-MiB. שימו לב ש-sharedMemPerMultiprocessor הוא המקסימום הניתן להקצאה (עד 228 KiB ב-Hopper), לא ברירת המחדל ל-block.

תרגיל 2 - מדידת latency של חישוב לעומת global memory

עכשיו נוכיח בפועל שגישה ל-global memory איטית בסדרי גודל מפעולת חישוב. נשתמש ב-clock64() שקורא את שעון ה-SM במחזורים.

  1. הוסיפו ל-memlat.cu שני kernels שרצים ב-thread בודד (<<<1,1>>>):
  2. fmachain - שרשרת FMA תלויה: a = a * c + d; שחוזרת STEPS פעמים, כשכל איטרציה תלויה בקודמת. מדדו את מספר המחזורים עם clock64() לפני ואחרי, וחלקו ב-STEPS.
  3. chase - "מרדף מצביעים" בזיכרון גלובלי: idx = buf[idx]; שחוזר STEPS פעמים, כשכל טעינה תלויה בערך שנטען קודם. מדדו באותה דרך.
  4. בנו את buf על ה-host כמערך של int בגודל N גדול משמעותית מ-L2 (למשל 64M ints = 256 MB > 50 MB), כך ש-buf[i] = (i + STRIDE) % N עם STRIDE גדול וזר ל-N (למשל מספר ראשוני). זה יוצר מעגל ארוך אחד של גישות מפוזרות שמחטיאות את הcache.
  5. העתיקו את buf ל-device, הריצו כל kernel, והדפיסו מחזורים לפעולה עבור שניהם.
  6. חשבו את היחס: פי כמה טעינה תלויה מ-global memory איטית מ-FMA תלויה?

רמז: כדי שהcompiler לא ימחק את הלולאות, כל איטרציה חייבת להיות תלויה בקודמת, ותוצאת הסוף חייבת להיכתב ל-global memory (*sink = ...). הריצו ב-<<<1,1>>> כי אתם מודדים latency של פעולה בודדת (שרשרת תלויה), לא throughput. clock64() מחזיר long long של מחזורי SM.

תרגיל 3 - איפה חי אוגר שגלש, וכמה הוא עולה

בשיעור אמרנו שאוגרים עודפים גולשים ל-local memory שיושב ב-GPU RAM. עכשיו נאלץ גלישה ונראה אותה בעיניים.

  1. כתבו kernel שמשתמש בהרבה אוגרים - למשל לולאה שמחזיקה מערך מקומי גדול (float acc[64];) שכולו נגזר מהקלט ומצטבר לתוצאה אחת.
  2. קמפלו פעם אחת עם nvcc -arch=sm_90 -Xptxas -v והסתכלו בשורת ה-ptxas info: כמה אוגרים נוצלו? האם יש spill stores / spill loads?
  3. עכשיו הגבילו בכוונה את האוגרים עם -maxrregcount=32 (או -Xptxas -maxrregcount=32) וקמפלו שוב. השוו: כמה אוגרים נוצלו עכשיו, וכמה בייטים של spill הופיעו?
  4. שלפו את ה-SASS עם cuobjdump -sass וחפשו הוראות STL (Store Local) ו-LDL (Load Local). הסבירו: להיכן פיזית ניגשות ההוראות האלה, ולמה הן יקרות פי כ-100 מגישה לאוגר?

רמז: STL/LDL הן גישות ל-local memory, שפיזית יושב ב-GPU RAM (DRAM) - לכן העלות דומה ל-global memory, מאות מחזורים. -maxrregcount מכריח את ptxas להשתמש בפחות אוגרים, ומה שלא נכנס גולש. השוו את שורות ה-spill stores/spill loads בין שתי הcompilations.

תרגיל 4 - "cache מנוהל-מתכנת" מול L1 של CPU

זהו תרגיל הסבר וניסוי קצר שממחיש את ההבדל התפיסתי הגדול של השיעור.

  1. כתבו בפסקה קצרה, במילים שלכם: מה ההבדל בין cache מנוהל-חומרה (L1 של CPU) לבין מנוהל-מתכנת (shared memory של GPU)? מי מחליט מה יושב בcache בכל אחד מהמקרים?
  2. כתבו שני kernels שמסכמים מערך קטן (למשל 256 float) 1000 פעמים: הראשון קורא בכל פעם מ-global memory ישירות; השני טוען פעם אחת ל-__shared__ float tile[256]; (עם __syncthreads()), ואז מסכם 1000 פעמים מה-shared memory.
  3. מדדו את שניהם (עם אירועי CUDA או clock64()) והשוו. איזה מהיר יותר ולמה?
  4. ענו: מדוע ב-CPU לא צריך לכתוב את הגרסה השנייה ידנית, אבל ב-GPU כן? קשרו לתשובה שלכם מסעיף 1.

רמז: ב-CPU החומרה תשמור אוטומטית את 256 האיברים ב-L1 אחרי הקריאה הראשונה, ו-999 הקריאות הבאות יהיו hit-ים ללא קוד מיוחד. ב-GPU אין הבטחה כזו ל-L1 - אתם חייבים לבצע את ה-staging ל-shared memory ביד כדי לקבל את אותו יתרון. זו בדיוק המשמעות של "מנוהל-מתכנת".

תרגיל 5 - לחץ אוגרים ותפוסה

נחבר בין קובץ האוגרים לתפוסה במספרים.

  1. עבור ה-H100 (65,536 אוגרים ל-SM, עד 2048 threads דרים), מלאו את הטבלה: לכל ערך של אוגרים-ל-thread מתוך {32, 40, 64, 96, 128, 255}, כמה threads מקסימלי יכולים לדור, וכמה זו תפוסה באחוזים?
  2. עבור ה-kernel מתרגיל 3, קחו את מספר האוגרים שדווח (לפני ההגבלה) והעריכו את התפוסה המקסימלית שהוא מאפשר.
  3. השתמשו ב-cudaOccupancyMaxActiveBlocksPerMultiprocessor (או במחשבון התפוסה) כדי לאמת את ההערכה שלכם עבור גודל block נתון.
  4. הסבירו: באיזה מקרה כדאי להוסיף -maxrregcount נמוך יותר, ומתי זה יזיק במקום לעזור?

רמז: threads מקסימלי = floor(65536 / regs_per_thread), מוגבל למקסימום 2048; תפוסה = threads/2048. הורדת אוגרים מעלה תפוסה אבל אם היא גורמת לגלישה - הרווח מהתפוסה עלול להיבלע על ידי מחיר ה-spills. זו פשרה, לא כלל.

תרגיל 6 (בונוס) - צוק ה-latency על פני השכבות

עד כה מדדנו קצה אחד (HBM) מול חישוב. עכשיו נצייר את כל העקומה.

  1. קחו את kernel ה-chase מתרגיל 2 והריצו אותו על מערכים בגדלים עולים: 16 KiB, 64 KiB, 256 KiB, 1 MiB, 4 MiB, 32 MiB, 128 MiB, 512 MiB (שמרו על מרדף מצביעים תלוי בכל גודל).
  2. עבור כל גודל, רשמו מחזורים-לגישה, וציירו גרף של latency מול גודל מערך העבודה (working set).
  3. זהו את "המדרגות" בגרף: היכן ה-working set נכנס ל-L1, היכן הוא עולה על L1 ונופל ל-L2, והיכן הוא עולה על L2 (50 MB) ונופל ל-HBM. השוו את המדרגות לגדלים מהפירמידה של תרגיל 1.

רמז: כל עוד ה-working set נכנס לשכבה, ה-latency נמוך; ברגע שהוא חורג, כל גישה מחטיאה ונופלת לשכבה הבאה, וה-latency קופץ. תראו בערך שלוש רמות: L1 (עשרות מחזורים), L2 (מאות בודדות), HBM (מאות רבות). הצוק בין 50 MB ל-64 MB הוא הגבול של ה-L2 של H100.