לדלג לתוכן

2.1 Threads ו Warps מודל SIMT פתרון

פתרון - Threads ו-Warps - מודל SIMT

נעבור תרגיל אחרי תרגיל. כל הפלטים כאן הם מ-H100 (SXM) עם CUDA 12.4; זמני הריצה המדויקים בתרגילי המדידה (4, 5, 6) ישתנו אצלכם בהתאם לכרטיס, לגרסת ה-toolkit ולעומס, וזה תקין - מה שקבוע הוא היחסים ולא המספרים המוחלטים. גודל ה-warp הוא 32 בכל הכרטיסים הרלוונטיים, ולכן כל החישובים של תרגילים 1-3 מדויקים גם על T4.

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

הטבלה המלאה (warpsPerBlock = ceil(size/32), פעילים באחרון = size - 32*(warps-1)):

 block size | warps | active in last warp | wasted lanes
     32     |   1   |         32           |        0
     64     |   2   |         32           |        0
    128     |   4   |         32           |        0
    256     |   8   |         32           |        0
    512     |  16   |         32           |        0
   1024     |  32   |         32           |        0
     96     |   3   |         32           |        0
    100     |   4   |          4           |       28
     33     |   2   |          1           |       31

חישוב ה-warp rank ל-block דו-ממדי (16,16) (מיפוי row-major, linearIdx = x + y*16):

thread (x=8, y=2):  linearIdx = 8 + 2*16 = 40  -> warpRank = 40/32 = 1,  lane = 40%32 = 8
thread (x=0, y=1):  linearIdx = 0 + 1*16 = 16  -> warpRank = 16/32 = 0,  lane = 16%32 = 16

הרצת ה-kernel עם block של 100 threads:

$ nvcc -arch=sm_90 warps.cu -o warps
$ ./warps 1
block 0 | warpRank 0/4 | first linearIdx 0
block 0 | warpRank 1/4 | first linearIdx 32
block 0 | warpRank 2/4 | first linearIdx 64
block 0 | warpRank 3/4 | first linearIdx 96

הודפסו בדיוק 4 שורות = 4 warps, תואם ל-ceil(100/32)=4. שימו לב ל-warp השלישי (rank 3) שמתחיל ב-linearIdx 96 ומכיל רק 4 threads פעילים (96,97,98,99) - 28 נתיבים ריקים. אחרי השינוי ל-dim3 block(16,16):

$ ./warps 1
block 0 | warpRank 0/8 | first linearIdx 0
block 0 | warpRank 1/8 | first linearIdx 32
...
block 0 | warpRank 7/8 | first linearIdx 224

ה-thread (8,2) עם linearIdx 40 נופל בטווח [32,63] של warp rank 1 - בדיוק כפי שחישבנו ביד.

למה זה עבד: מספר ה-warps הוא עיגול כלפי מעלה כי החומרה חייבת להקצות warp שלם גם עבור thread בודד; ה-warp rank הוא פשוט האינדקס הליניארי חלקי 32. איך להכליל: תמיד לינאריזציה של ה-thread לפי row-major (x הכי מהיר, z הכי איטי) ואז חלוקה ב-warpSize - זה קובע איזה thread שכן לאיזה, ומכאן נגזרים coalescing ו-divergence.

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

הרצה עם warp מלא (32 threads):

$ ./warps 2
warp @ block 0: activemask=0xffffffff  evens=0x55555555
  • activemask=0xffffffff - כל 32 הביטים דלוקים, כי כל 32 הנתיבים פעילים. __popc(0xffffffff)=32.
  • evens=0x55555555 - הנתיבים הזוגיים (0,2,4,...,30) תרמו 1. בבינארי זו התבנית 0101...0101, שהיא 0x5 בכל ניבל, כלומר 0x55555555. __popc(evens)=16 (16 נתיבים זוגיים מתוך 32).

אחרי השינוי ל-showMask<<<1, 20>>>():

$ ./warps 2
warp @ block 0: activemask=0x000fffff  evens=0x00055555

עכשיו רק 20 הביטים הנמוכים דלוקים ב-activemask (0x000fffff, כלומר lanes 0..19) - זהו warp חלקי: ה-block מכיל רק 20 threads, אז ה-warp היחיד רץ עם 20 נתיבים פעילים ו-12 מכובים. ה-evens מכיל כעת רק את הזוגיים 0,2,...,18 = 10 ביטים = 0x00055555. עם השורה שמוסיפה __popc:

warp @ block 0: activemask=0x000fffff (popc=20)  evens=0x00055555 (popc=10)

למה זה עבד: __activemask() הוא צילום של הנתיבים הפעילים ברגע הקריאה, ו-__ballot_sync בונה מספר שבו ביט i מייצג את ה-predicate של נתיב i - כך ה-warp הופך לישות שאפשר לבחון כמסכת ביטים. איך להכליל: מסכות הן שפת התקשורת של warp-level primitives; __popc על מסכה סופר משתתפים, __ffs מוצא את הנתיב הפעיל הראשון, ומעביר את המסכה כארגומנט של פרימיטיב _sync מבטיח שכל המשתתפים ייפגשו באותה הוראה.

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

הקוד המלא שנוסף ל-warps.cu:

__inline__ __device__ float warpReduceSum(float val) {
    for (int offset = warpSize / 2; offset > 0; offset >>= 1)
        val += __shfl_down_sync(0xffffffff, val, offset);
    return val;
}
__global__ void reduceKernel() {
    float v  = threadIdx.x;                 // values 0,1,...,blockDim-1
    float s  = warpReduceSum(v);
    int   lane = threadIdx.x % warpSize;
    if (lane == 0)
        printf("warp starting at tid %d: sum = %.0f\n", threadIdx.x, s);
}

הרצה עם block של 32:

$ ./warps 3
warp starting at tid 0: sum = 496

הסכום 0+1+...+31 = (3132)/2 = 496*, בדיוק. עם אתחול v = 1.0f במקום threadIdx.x:

warp starting at tid 0: sum = 32

הסכום הוא 32 - כלומר הצמצום ספר את מספר הנתיבים הפעילים ב-warp. הרצה עם block של 64 (שני warps):

$ ./warps 3
warp starting at tid 0:  sum = 496
warp starting at tid 32: sum = 1520

warp 0 מסכם את הערכים 0..31 = 496; warp 1 מסכם את 32..63 = (32+63)32/2 = 1520. כל warp מצטמצם בנפרד* כי __shfl_down_sync לעולם אינו חוצה גבול warp. כדי לקבל סכום של כל ה-block צריך לצבור את התוצאות של שני ה-warps דרך shared memory - כל warp כותב את הסכום החלקי שלו ל-shared, ואז warp אחד מסכם את החלקיים (נלמד את התבנית בפרק 3).

למה זה עבד: חמשת שלבי ה-shuffle (16,8,4,2,1) מקפלים את ה-warp לחצי בכל שלב, ו-log2(32)=5 שלבים מספיקים לצמצם 32 ערכים לאחד, כשכל התקשורת דרך האוגרים בלבד. איך להכליל: reduction ברמת warp הוא אבן היסוד של כל צמצום מהיר (סכום, max, dot product); הוא מהיר מ-shared memory כי הנתונים לא עוזבים את קובץ האוגרים, ומרכיבים ממנו צמצום ברמת block על ידי שילוב עם shared memory בין ה-warps.

פתרון תרגיל 4 - נכון תמיד, מהיר לפעמים

שני ה-kernels מייצרים תוצאות נכונות: copyCoalesced נותן out == in, ו-copyStrided נותן out[t] == in[t*stride % n]. שני הפלטים נכונים כשלעצמם - זו בדיוק הנקודה. המדידה (n = 2^24 floats, block=256, stride=1024):

$ ./warps 4
copyCoalesced :   48.6 us   (~2.6 TB/s effective)
copyStrided   :  372.1 us   (~0.34 TB/s effective)
ratio         :  ~7.7x slower

התוצאות זהות בנכונותן, אבל הגישה המפוזרת איטית פי ~8. ההסבר במונחי warp:

coalesced warp:            32 lanes -> 32 contiguous addresses
                          in[i], in[i+1], ..., in[i+31]
                          = 128 contiguous bytes -> one transaction of 128 bytes (4 sectors)

strided warp:              32 lanes -> 32 addresses spread apart by stride*4 bytes
                          each one falls in a different sector
                          -> up to 32 separate transactions for the same 32 threads

אותה עבודה לוגית (32 קריאות של 4 בתים), אבל פי 8 תעבורת זיכרון בפועל, כי החומרה נאלצת למשוך מקטע של 32 בתים עבור כל נתיב במקום מקטע אחד ל-32 נתיבים.

למה זה עבד: המנוע שמאחד גישות (coalescing) פועל ברמת ה-warp - הוא בוחן את 32 הכתובות של ה-warp בבת אחת וממזג אותן לכמה שפחות עסקאות; רק כשהן רצופות המיזוג מלא. איך להכליל: זו ההדגמה המזוקקת של "warps לא רלוונטיים לנכונות, חיוניים לביצועים" - סדרו תמיד את הקוד כך שנתיבים סמוכים ב-warp יגשו לכתובות סמוכות בזיכרון (in[blockIdx.x*blockDim.x + threadIdx.x]), ותקבלו coalescing בחינם; נעמיק בזה בפרק 3.

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

המדידה של שני המצבים (block=256, 2^20 threads):

$ ./warps 5
mode 0 (all 100 iters)          :   61 us
mode 1 (half 10, half 200)      :  118 us

ב-mode 1 ה-warp "משלם" על ~200 איטרציות, לא על הממוצע 105. הסיבה: ל-warp יש PC אחד וזרם הוראות אחד. כשחצי הנתיבים לוקחים את ענף ה-10 והחצי את ענף ה-200, החומרה מריצה את שני הענפים בזה אחר זה - קודם הנתיבים של ענף אחד פעילים והשאר מכובים (predicated off), ואז להפך. ה-warp אינו משתחרר עד שגם המסלול האיטי ביותר (200 איטרציות) הסתיים, ולכן זמן ה-warp נקבע על ידי הענף הארוך. וזה בדיוק מה שרואים: mode 1 איטי פי ~2 מ-mode 0 (200 מול 100), למרות שכמות העבודה ה"ממוצעת" קטנה יותר.

שינוי התנאי ל-(lane % 2 == 0):

mode 1 with (lane % 2 == 0)     :  ~118 us   (roughly the same)

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

למה זה עבד: מקור העונש הוא ה-PC היחיד של ה-warp; פיצול בקרה בתוך warp מסריאליזציה את הענפים. איך להכליל: כדי להימנע מ-divergence, ארגנו את הבקרה כך שכל 32 הנתיבים של warp ילכו לאותו מסלול - למשל שכל warp שלם יטפל בסוג נתונים אחד, במקום להחליט לפי הנתיב. הטיפול המלא, כולל independent thread scheduling של Volta ומדידת ה-divergence ב-profiler, בשיעור 8.3.

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

פירוק ה-SASS של copyCoalesced על sm_90:

$ cuobjdump -sass warps | sed -n '/copyCoalesced/,/copyStrided/p'
      /*0000*/  MOV R1, c[0x0][0x28] ;
      /*0010*/  S2R R0, SR_CTAID.X ;        // blockIdx.x  (the block id)
      /*0020*/  S2R R3, SR_TID.X ;          // threadIdx.x (the thread id within the warp)
      /*0030*/  IMAD R0, R0, c[0x0][0x0], R3 ;   // i = blockIdx.x*blockDim.x + threadIdx.x
      /*0040*/  ISETP.GE.AND P0, PT, R0, c[0x0][0x180], PT ;  // i < n ?
      /*0050*/  ...
      /*0070*/  LDG.E R4, [R4.64] ;         // load from global: in[i]
      /*0090*/  STG.E [R6.64], R4 ;         // store to global: out[i]

ההוראה S2R R3, SR_TID.X היא בדיוק נקודת המפגש: מודל התכנות הסקלרי (threadIdx.x) הופך לאוגר פר-thread שממנו כל אחד מ-32 נתיבי ה-warp מקבל ערך שונה. הרצת שני ה-kernels תחת ncu עם מדד ה-sectors:

$ ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum \
      --launch-count 2 ./warps 4
  copyCoalesced   l1tex__t_sectors...ld.sum   sector   2,097,152
  copyStrided     l1tex__t_sectors...ld.sum   sector  16,777,216

אימות החשבון: יש 2^24 / 32 = 524,288 warps. גישה מאוחדת מושכת 4 sectors לכל warp (128 בתים / 32) -> 524,288 x 4 = 2,097,152. גישה מפוזרת מושכת עד 32 sectors לכל warp (sector לכל נתיב) -> 524,288 x 32 = 16,777,216 - פי 8 בדיוק, כמו יחס הזמנים בתרגיל 4.

למה זה עבד: מדד ה-sectors סופר כמה מקטעי 32 בתים נמשכו בפועל מ-L2, וזה מכמת ישירות את איכות ה-coalescing של ה-warp; היחס פי 8 בין מפוזר למאוחד הוא בדיוק ההסבר למרווח הזמנים. איך להכליל: cuobjdump -sass/nvdisasm חושפים איך מודל התכנות מתורגם להוראות פר-warp (S2R, LDG, STG), ו-ncu מודד את המחיר; כשמכווננים ביצועים, מדד ה-sectors לכל בקשה הוא המדד הישיר ל"עד כמה ה-warps שלי מאוחדים" - נשתמש בו שוב בפרק 3.