לדלג לתוכן

3.1 היררכיית הזיכרון הלוגית תרגול

תרגול - היררכיית הזיכרון הלוגית

בתרגול הזה תמפו בעצמכם את היררכיית הזיכרון על קוד אמיתי. תסמנו היכן חי כל משתנה ב-kernel נתון, תגרמו למשתנה לגלוש ל-local memory ותראו את הגלישה בעיניים בעזרת ה-ptxas, תוכיחו במדידה ש-shared memory הוא פרטי ל-block, ולבסוף תנתחו איזה תיאום בין blocks חוקי ואיזה מוביל ל-deadlock. חלק מהתרגילים הם ניתוחיים (קריאת פלט compiler, נימוק על נכונות) וחלק דורשים הרצה על GPU. עבדו לפי הסדר. הריצו על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אך כל התרגילים ירוצו על כל GPU של NVIDIA.

הכנה

צרו קובץ mem.cu. לאורך התרגול נשתמש בדגל -v של ה-ptxas כדי לראות כמה registers ו-local memory כל kernel צורך - זה הכלי המרכזי שלנו לצפייה בהיררכיית הזיכרון מבפנים:

nvcc -O2 -arch=sm_90a --ptxas-options=-v -o mem mem.cu

החליפו את sm_90a בארכיטקטורה של הכרטיס שלכם (למשל sm_75 ל-T4). לבירור ה-compute capability:

nvidia-smi --query-gpu=name,compute_cap --format=csv

תרגיל 1 - מיפוי מרחבי הזיכרון (ניתוח)

לפניכם kernel. בלי להריץ דבר, סַמנו עבור כל משתנה שמסומן באיזה מרחב זיכרון הוא חי: register, local, shared, global, או constant.

__constant__ float coeff[4];                          // (A)

__global__ void analyze(const float* in, float* out, int n) {
    __shared__ float tile[128];                       // (B)
    int   i   = blockIdx.x * blockDim.x + threadIdx.x;// (C)
    float acc = 0.0f;                                 // (D)
    float buf[8];                                     // (E)
    tile[threadIdx.x] = in[i];                        // in -> (F)
    __syncthreads();
    for (int k = 0; k < 4; ++k)
        acc += coeff[k] * tile[threadIdx.x];          // (G)
    for (int k = 0; k < 8; ++k)
        buf[k % n] = acc;                             // (H)
    out[i] = buf[i % 8];                              // out -> (I)
}
  1. עבור כל אחד מהסימונים A עד I, כתבו את מרחב הזיכרון.
  2. שני משתנים חולקים scope זהה (thread בודד) אבל מיקום פיזי מנוגד - אילו, ומדוע נבדלים.
  3. הסבירו: coeff ו-in שניהם בהיקף grid; מה ההבדל ביניהם?

רמז: משתנה סקלרי מקומי יושב ב-register. מערך מקומי שניגשים אליו באינדקס לא-קבוע בזמן קומפילציה (שימו לב ל-buf[k % n] ול-buf[i % 8]) לא יכול לשבת באוגרים. __shared__ תמיד shared, __constant__ תמיד constant, ומצביע שהתקבל מ-cudaMalloc מצביע ל-global.

תרגיל 2 - לגרום לגלישה ל-local memory ולראות אותה

  1. כתבו שני kernels בקובץ mem.cu:
__global__ void tight(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i] * 2.0f + 1.0f;
}

__global__ void heavy(const int* idx, const float* in, float* out, int n) {
    float scratch[64];                       // large; a candidate to spill
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    for (int k = 0; k < 64; ++k) scratch[k] = in[(i + k) % n];
    float s = 0.0f;
    for (int k = 0; k < 64; ++k) s += scratch[idx[i] % 64] * scratch[k];
    out[i] = s;
}
  1. הוסיפו main מינימלי שרק משיק את שני ה-kernels (לא חייב תוצאה נכונה - אנחנו בודקים קומפילציה בלבד).
  2. קמפלו עם --ptxas-options=-v וקראו את פלט ה-ptxas לכל kernel.
  3. רשמו לכל kernel: כמה registers וכמה stack frame / spill stores / spill loads / lmem הוא צורך. איזה מהם גלש ל-local memory?

רמז: חפשו בפלט שורות כמו Used NN registers ו-NN bytes stack frame, NN bytes spill stores, NN bytes spill loads. ל-tight תראו 0 בתים stack frame; ל-heavy, בגלל המערך scratch[64] באינדוקס דינמי, תראו stack frame גדול מ-0 - זהו ה-local memory. ככל שתגדילו את scratch (נסו 128 או 256), הגלישה תגדל.

תרגיל 3 - shared memory הוא פרטי ל-block (הוכחה במדידה)

נוכיח ש-shared memory של block אחד אינו נראה ל-block אחר.

  1. כתבו kernel שבו כל block כותב את מספרו לתוך התא הראשון של ה-shared שלו, ואז כל thread קורא משם:
__global__ void privacy(int* out) {
    __shared__ int marker;
    if (threadIdx.x == 0) marker = blockIdx.x;   // every block writes its own identity
    __syncthreads();
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    out[gid] = marker;                            // every thread reports what it saw
}
  1. השיקו עם blockSize = 4 ו-numBlocks = 3 (12 threads בסך הכל), העתיקו את out חזרה והדפיסו את 12 הערכים.
  2. בדקו: האם threads ב-block 2 ראו אי-פעם את הערך של block 0 או block 1?
  3. הסבירו מה זה מוכיח על ההיקף של __shared__.

רמז: אם ה-shared היה משותף בין blocks, הייתם רואים ערך אחיד (כולם רואים את מי שכתב אחרון). במקום זאת תראו שכל רביעיית threads מדווחת את blockIdx.x של ה-block שלה עצמו - הוכחה ש-marker הוא עותק נפרד לכל block.

תרגיל 4 - תיאום חוקי מול deadlock בין blocks (ניתוח)

לפניכם שני kernels שמנסים לתאם בין blocks. עבור כל אחד, הכריעו: חוקי ובטוח, או עלול ל-deadlock / להיות שגוי - ונמקו.

// version A: every block adds to a shared global counter
__global__ void versionA(const float* in, int* counter, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n && in[i] > 0.0f) atomicAdd(counter, 1);
}

// version B: block 0 waits until the last block signals it's done
__global__ void versionB(volatile int* flag) {
    if (blockIdx.x == gridDim.x - 1 && threadIdx.x == 0)
        *flag = 1;                       // the last block signals
    if (blockIdx.x == 0 && threadIdx.x == 0)
        while (*flag == 0) { /* spin */ }// block 0 spin-waits
}
  1. עבור גרסה A: האם היא נכונה? מה מבטיח את הנכונות?
  2. עבור גרסה B: תארו תרחיש קונקרטי שבו היא נתקעת ל-deadlock.
  3. הציעו את הדרך הנכונה להשיג את הסנכרון הגלובלי שגרסה B ניסתה להשיג.

רמז: זכרו מההרצאה שסדר הביצוע בין blocks אינו דטרמיניסטי, ולא כל ה-blocks בהכרח תושבים בו-זמנית. אם ל-GPU יש פחות "מקומות" מ-blocks, ה-block האחרון עלול לא להתחיל לרוץ עד ש-block 0 יפַנה מקום - אבל block 0 תקוע בלולאת ההמתנה. atomicAdd, לעומת זאת, לא מניח שום סדר. הסנכרון הגלובלי הנכון: לפצל לשני kernels.

תרגיל 5 - registers מול occupancy (ניתוח כמותי)

השתמשו בפלט ה-registers מתרגיל 2 כדי לחשב את השפעת ה-registers על התפוסה.

  1. הניחו ש-heavy צורך R registers ל-thread (קחו את המספר האמיתי מהפלט שלכם). ל-SM של H100 יש 65,536 אוגרים בני 32 סיביות.
  2. חשבו את המספר המרבי של threads תושבים שקובץ האוגרים מאפשר: 65536 / R.
  3. השוו למקסימום הפיזי של ה-H100 - 2,048 threads תושבים לכל SM. האם קובץ האוגרים הוא הגורם המגביל?
  4. חזרו על החישוב עבור ה-kernel הקל tight וראו את ההבדל.

רמז: אם 65536 / R קטן מ-2,048, אז ה-registers מגבילים את התפוסה - יש לכם פחות warps תושבים, ולכן פחות יכולת להסתיר latency. אם 65536 / R גדול מ-2,048, הגבול הוא בכלל אחר (מספר ה-threads המקסימלי), וה-registers לא מכבידים. שימו לב שהקצאה נעשית בפועל בגרנולריות של warp; החישוב כאן הוא קירוב לחסם העליון.

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

  1. הפיקו PTX ל-kernel tight מתרגיל 2:
nvcc -arch=sm_90a -ptx -o tight.ptx mem.cu
  1. פתחו את tight.ptx ומצאו את גוף ה-kernel tight. אתרו את הצהרות ה-.reg (הכרזות על רגיסטרים וירטואליים) ואת פקודות ה-ld.global / st.global (קריאה וכתיבה ל-global memory).
  2. זהו: אילו רגיסטרים וירטואליים מסוג .f32 (נקודה צפה 32 סיביות) משמשים לחישוב in[i] * 2.0f + 1.0f, ואילו מסוג .u32 / .s32 משמשים לחישוב האינדקס i?
  3. חשבו: כמה ld.global יש בגוף, וכמה st.global? האם זה תואם לכך שהחישוב memory-bound?

רמז: ה-PTX חושף בדיוק את מה שהיה בלתי-נראה ב-CUDA C - הרגיסטרים. חפשו שורות כמו .reg .f32 %f<N>; ו-.reg .b32 %r<N>;. פקודת ld.global.f32 היא הקריאה מ-global, st.global.f32 היא הכתיבה. תראו קריאה אחת (in[i]) וכתיבה אחת (out[i]) - יחס גישות זיכרון גבוה לפעולה אריתמטית בודדת, בדיוק הסימן לחישוב memory-bound.