לדלג לתוכן

2.1 Threads ו Warps מודל SIMT תרגול

תרגול - Threads ו-Warps - מודל SIMT

בתרגול הזה נהפוך את ה-warp מרעיון מופשט לדבר שרואים ומודדים. נחשב ידנית כמה warps יש ב-block בגדלים שונים ומה ה-warp rank של thread נתון, ואז נאמת את החישוב על GPU אמיתי; נשתמש ב-__activemask() וב-__ballot_sync() כדי "לראות" את 32 הנתיבים של warp; נבנה reduction ברמת warp עם __shfl_down_sync(); ולבסוף נראה את הנקודה החשובה ביותר של השיעור בעיניים - kernel שהתוצאה שלו נכונה ללא תלות ב-warps, אבל המהירות שלו תלויה בהם לחלוטין (coalescing מול גישה מפוזרת). עבדו לפי הסדר. אפשר לעבוד על כרטיס NVIDIA מקומי או ב-Google Colab (T4 חינם). אם אתם ב-Colab, החליפו -arch=sm_90 ב--arch=sm_75 והוסיפו ! לפני כל פקודת shell. גודל ה-warp הוא 32 בכל הכרטיסים האלה, כך שכל החישובים תקפים; רק תקרת ה-warps ל-SM ומספר ה-SM שונים (ל-T4: 40 SM, 32 warps ל-SM).

הכנה - קובץ מקור לחקירת warps

צרו את warps.cu. הוא יכיל כמה kernels שנפעיל בתרגילים השונים, ומקרו לבדיקת שגיאות שנשתמש בו לכל אורך הקורס:

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

// exercise 1: every thread prints its warp rank and lane
__global__ void reportWarps() {
    int linearIdx = threadIdx.x
                  + threadIdx.y * blockDim.x
                  + threadIdx.z * blockDim.x * blockDim.y;
    int warpRank = linearIdx / warpSize;
    int lane     = linearIdx % warpSize;
    int warpsPerBlock = (blockDim.x * blockDim.y * blockDim.z + warpSize - 1) / warpSize;
    // print only lane 0 of each warp so we don't flood the output
    if (lane == 0) {
        printf("block %d | warpRank %d/%d | first linearIdx %d\n",
               blockIdx.x, warpRank, warpsPerBlock, linearIdx);
    }
}

// exercise 2: showing the mask of the 32 lanes
__global__ void showMask() {
    unsigned active = __activemask();
    // every lane whose threadIdx.x is even contributes 1 to the vote
    unsigned evens  = __ballot_sync(active, (threadIdx.x % 2) == 0);
    int lane = threadIdx.x % warpSize;
    if (lane == 0) {
        printf("warp @ block %d: activemask=0x%08x  evens=0x%08x\n",
               blockIdx.x, active, evens);
    }
}

int main(int argc, char** argv) {
    int which = (argc > 1) ? atoi(argv[1]) : 1;
    if (which == 1) {
        dim3 block(100);            // intentionally not a multiple of 32
        reportWarps<<<1, block>>>();
    } else if (which == 2) {
        showMask<<<1, 32>>>();      // a single warp
    }
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    return 0;
}

קמפלו עם: nvcc -arch=sm_90 warps.cu -o warps

תרגיל 1 - warpsPerBlock ו-warp rank ביד ובקוד

נתרגל את הנוסחאות מהשיעור, ואז נאמת אותן.

  1. בלי מחשב, מלאו טבלה: עבור גודלי block של 32, 64, 128, 256, 512, 1024, 96, 100, 33 - חשבו warpsPerBlock = ceil(size/32), כמה threads פעילים יש ב-warp האחרון, וכמה נתיבים מבוזבזים.
  2. עבור block דו-ממדי בגודל (16, 16), חשבו ביד את ה-linearIdx, ה-warpRank וה-lane של ה-thread בקואורדינטות (threadIdx.x=8, threadIdx.y=2) ושל (threadIdx.x=0, threadIdx.y=1).
  3. הריצו ./warps 1 (ה-block הוא 100 threads). ספרו כמה שורות הודפסו - זה בדיוק מספר ה-warps ב-block. האם הוא תואם ל-ceil(100/32) שחישבתם?
  4. שנו את גודל ה-block ב-main מ-dim3 block(100) ל-dim3 block(16,16), קמפלו והריצו שוב. אתרו בפלט את ה-warp שאליו שייך ה-thread (8,2) ואשרו מול החישוב הידני שלכם.

רמז: מספר ה-warps ל-block הוא תמיד עיגול כלפי מעלה, לא כלפי מטה - block של 100 הוא 4 warps, לא 3. ב-block דו-ממדי המיפוי הוא row-major: linearIdx = threadIdx.x + threadIdx.y * blockDim.x, ולכן ב-block ברוחב 16, כל warp מכסה שתי שורות שלמות (16 + 16 = 32).

תרגיל 2 - לראות את 32 הנתיבים עם activemask ו-ballot

עכשיו נסתכל ישירות על ה-warp כמסכת ביטים.

  1. הריצו ./warps 2. רשמו את הערכים של activemask ושל evens. כמה ביטים דלוקים ב-activemask (השתמשו במחשבון או ב-__popc מנטלית)? למה בדיוק כך?
  2. הסבירו את הערך של evens: אילו נתיבים תרמו 1? כתבו את ה-16 הביטים הנכונים בבינארי וודאו שהם תואמים לפלט ההקסה.
  3. שנו את ה-launch ב-main מ-showMask<<<1, 32>>>() ל-showMask<<<1, 20>>>() (block קטן מ-warp). הריצו שוב. כמה ביטים דלוקים כעת ב-activemask? מה זה מלמד על "warp חלקי"?
  4. הוסיפו ל-showMask שורה שמדפיסה גם את __popc(active) (ספירת הביטים הדלוקים) ואת __popc(evens). הריצו שוב עם block של 32 ושל 20 ואשרו את הספירות.

רמז: __activemask() מחזיר ביט לכל נתיב פעיל בזמן הקריאה. עבור warp מלא של 32 זה 0xffffffff (32 ביטים). __ballot_sync(mask, pred) בונה מספר שבו הביט ב-lane i דלוק אם ה-predicate של lane i אמת - כאן, הנתיבים הזוגיים, כלומר התבנית ...01010101 = 0x55555555. עם 20 נתיבים בלבד, רק 20 הביטים הנמוכים יכולים להידלק.

תרגיל 3 - reduction ברמת warp עם shfl_down_sync

נבנה את הצמצום ברמת ה-warp מהשיעור ונאמת שהוא נכון.

  1. הוסיפו ל-warps.cu את הפונקציה warpReduceSum מהשיעור (עם הלולאה for (int offset = warpSize/2; offset > 0; offset >>= 1)).
  2. כתבו kernel reduceKernel שבו כל thread מאתחל float v = threadIdx.x; (כלומר הערכים 0,1,2,...,31), קורא ל-warpReduceSum(v), ונתיב 0 מדפיס את התוצאה.
  3. הריצו עם block של 32. התוצאה הצפויה היא סכום 0+1+...+31. חשבו אותו ביד ((31*32)/2) ואשרו.
  4. שנו כך שכל thread מאתחל v = 1.0f והריצו שוב - כמה אתם מצפים לקבל, ולמה? זה סופר בעצם את מספר הנתיבים הפעילים.
  5. (אתגר) הריצו את reduceKernel עם block של 64. עכשיו יש שני warps. מה מדפיס נתיב 0 של כל warp? הסבירו למה כל warp מצטמצם בנפרד, ומה תצטרכו להוסיף כדי לקבל סכום של כל ה-block (רמז: shared memory, נלמד בפרק 3).

רמז: __shfl_down_sync(0xffffffff, val, offset) מחזיר את ה-val של הנתיב lane + offset (ואם הוא מחוץ ל-warp, ערך לא-מוגדר שלא נשתמש בו). חמשת השלבים 16,8,4,2,1 מקפלים את ה-warp לחצי בכל פעם, וה-log2(32)=5 שלבים מספיקים. שימו לב ש-__shfl_down_sync פועל בגבולות warp בלבד; הוא לעולם לא חוצה ל-warp אחר, ולכן block של 64 נותן שני סכומים נפרדים.

תרגיל 4 - נכון תמיד, מהיר לפעמים (coalescing מול גישה מפוזרת)

זה התרגיל המרכזי: kernel שהתוצאה שלו נכונה ללא קשר ל-warps, אבל המהירות שלו תלויה בהם. נעתיק וקטור בשתי דרכים.

  1. הוסיפו ל-warps.cu שני kernels שמעתיקים מערך float באורך n מ-in ל-out:
// coalesced access: thread i reads element i -> 32 lanes = 32 contiguous elements
__global__ void copyCoalesced(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i];
}

// strided access: the exact same computation, but a stride skip between adjacent lanes
__global__ void copyStrided(const float* in, float* out, int n, int stride) {
    int t = blockIdx.x * blockDim.x + threadIdx.x;
    int i = (long long)(t) * stride % n;   // scrambling the address each lane accesses
    if (t < n) out[t] = in[i];
}
  1. שכתבו את main כך שיקצה n = 1<<24 floats על ה-device (השתמשו ב-CUDA_CHECK סביב cudaMalloc), יאתחל את הקלט, ויריץ כל kernel עם block=256. מדדו זמן עם cudaEvent (או פשוט הריצו תחת nsys/ncu).
  2. הריצו את שני ה-kernels וודאו ששניהם מייצרים את אותה תוצאה נכונה (לגישה המאוחדת, out שווה ל-in; לגישה המפוזרת, out[t] = in[t*stride % n] - ודאו נקודתית). הנקודה: התוצאה של כל אחד נכונה כשלעצמה, ללא תלות ב-warps.
  3. השוו את הזמנים של copyCoalesced מול copyStrided עם stride גדול (למשל 1024). איזה מהיר יותר, ובכמה בערך?
  4. הסבירו במונחי warp: מה קורה כש-32 הנתיבים של warp ניגשים ל-32 כתובות רצופות (coalesced) לעומת 32 כתובות מפוזרות? כמה עסקאות זיכרון החומרה צריכה בכל מקרה?

רמז: זו ההדגמה החיה של "warps לא רלוונטיים לנכונות, חיוניים לביצועים". שתי הפונקציות מחזירות ערכים נכונים. אבל כשה-warp ניגש ל-32 אלמנטים רצופים, החומרה מאחדת (coalesces) אותם לעסקה אחת רחבה של 128 בתים; כשהם מפוזרים, היא נאלצת להנפיק עד 32 עסקאות נפרדות - פי כמה יותר תעבורת זיכרון עבור אותה עבודה. נלמד את המנגנון במלואו בפרק 3; כאן רק מודדים את ההשפעה.

תרגיל 5 - התפצלות ה-warp במספרים (warp divergence)

נראה את ההשפעה של divergence, ההצצה מהשיעור.

  1. כתבו kernel divergent שבו כל thread מבצע לולאה שאורכה תלוי ב-lane שלו, למשל:
__global__ void divergent(float* out, int mode) {
    int lane = threadIdx.x % warpSize;
    float x = 0.0f;
    int iters;
    if (mode == 0)      iters = 100;                 // everyone equal - no divergence
    else                iters = (lane < 16) ? 10 : 200; // half the warp differs from the other half
    for (int k = 0; k < iters; ++k) x = fmaf(x, 1.001f, 1.0f);
    out[blockIdx.x * blockDim.x + threadIdx.x] = x;
}
  1. הריצו את שני המצבים (mode 0 ו-mode 1) ומדדו זמן. ב-mode 0 כל ה-warp מריץ 100 איטרציות; ב-mode 1 חצי מריץ 10 וחצי 200.
  2. שאלה: ב-mode 1, כמה איטרציות בפועל ה-warp "משלם" עליהן - 200 או 105 (הממוצע)? הסבירו למה, לאור העובדה של-warp יש PC אחד וההוראות של שני הענפים רצות בזה אחר זה.
  3. שנו את התנאי ל-(lane % 2 == 0) במקום (lane < 16). האם ה-divergence גרוע יותר, פחות, או אותו דבר? הסבירו.

רמז: ל-warp יש זרם הוראות אחד. כשחצי מהנתיבים רוצים 200 איטרציות והחצי 10, החומרה מריצה את שני המסלולים ברצף כשהנתיבים הלא-רלוונטיים מכובים (predicated off), ולכן ה-warp מתעכב עד שגם הנתיב האיטי ביותר סיים - כלומר משלם על ~200 איטרציות, לא על הממוצע. חלוקה lane < 16 לעומת lane % 2 נותנת אותו יחס עבודה, אבל שתיהן מפצלות את אותו warp - מה שקובע הוא אם ה-32 הנתיבים באותו warp מסכימים, לא איך הם מסודרים.

תרגיל 6 (בונוס) - אימות ה-warp דרך SASS

נסתכל מתחת למכסה המנוע ונראה את ה-warp בקוד המכונה.

  1. קמפלו את copyCoalesced ל-SASS: cuobjdump -sass warps | less (או nvcc -arch=sm_90 -cubin warps.cu -o warps.cubin && nvdisasm -c warps.cubin).
  2. אתרו את ההוראה שקוראת את %tid.x (מופיעה כ-S2R Rx, SR_TID.X) ואת ההוראה שקוראת את מזהה ה-block (SR_CTAID.X). זהו את הוראת הטעינה (LDG) ואת האחסון (STG).
  3. הריצו את שני ה-kernels (copyCoalesced ו-copyStrided) תחת ncu עם מדד ה-coalescing: ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum ./warps. השוו כמה sectors (מקטעי 32 בתים) כל kernel קרא. הגישה המפוזרת אמורה לקרוא הרבה יותר sectors עבור אותו מספר threads.
  4. הסבירו כיצד המספר הזה מכמת בדיוק את מה שראיתם בתרגיל 4: אותה עבודה לוגית, פי כמה תעבורת זיכרון בגלל סידור ה-warp.

רמז: ב-SASS, ההוראה S2R (Special-register to Register) היא איך ה-thread קורא את threadIdx/blockIdx שלו - זו הנקודה שבה מודל התכנות הסקלרי פוגש את ה-warp החומרתי. מדד ה-sectors ב-ncu מודד כמה מקטעי 32 בתים נמשכו מ-L2; גישה מאוחדת של warp מושכת 4 sectors (128 בתים) לכל 32 threads, וגישה מפוזרת יכולה למשוך עד 32 sectors - פי 8 תעבורה. זו אותה תובנה של תרגיל 4, עכשיו כמדד חומרה קשיח.