2.1 Threads ו Warps מודל SIMT הרצאה
בפרק 1 בנינו את ה-GPU מלמטה למעלה: ראינו ב-1.1 את חלוקת השבב ל-GPC, TPC ו-SM, פירקנו ב-1.2 את ה-SM (מעבד רב-זרמי - Streaming Multiprocessor) ואת קובץ האוגרים הענק שלו, מנינו ב-1.3 את הexecution pipelines (CUDA Core, SFU, LSU), וב-1.4 ראינו את ה-warp scheduler (מתזמן ה-warp) שמנפיק warp אחד למחזור שעון ומסתיר latency בהחלפות זולות. עד כה דיברנו בשפת החומרה. הפרק הזה פותח את הצד השני של הגלוסרי - מודל התכנות של CUDA, שפת המתכנת - ותפקידו הראשון של השיעור הזה הוא לגשר: לקחת את החומרה מ-1.2 עד 1.4 ולמפות אותה אל שתי המילים היסודיות שבהן המתכנת מבטא מקביליות, ה-thread וה-warp. נראה שה-thread הוא האטום של מודל התכנות, שה-warp הוא קבוצת 32 threads שהחומרה מתזמנת יחד, ושמודל ה-SIMT (Single Instruction, Multiple Threads) הוא נקודת המפגש ביניהם. התובנה החשובה ביותר של השיעור, שנחזור אליה שוב ושוב, היא עדינה: ה-warp כלל אינו חלק ממודל התכנות הרשמי - הוא פרט מימוש של החומרה, בדיוק כמו שורת cache ב-CPU. אפשר לכתוב CUDA נכון בלי לחשוב על warps; אי אפשר לכתוב CUDA מהיר בלי להבין אותם.
ה-thread - האטום של ההיררכיה¶
ה-thread הוא, בלשון הגלוסרי, "יחידת התכנות הנמוכה ביותר ל-GPU, הבסיס והאטום של היררכיית ה-threads של מודל התכנות של CUDA". זו היחידה הקטנה והבלתי-מתחלקת שהמתכנת חושב עליה: כשאתם כותבים __global__ void kernel(...), אתם כותבים את הקוד שירוץ ב-thread בודד, ואז מריצים ממנו עותקים רבים. כל שאר ההיררכיה שנפגוש בשיעורים הבאים - block (בלוק) ב-2.2, grid (רשת) ב-2.2 - היא רק דרך לארגן קבוצות של ה-thread הזה. ה-thread נמצא בתחתית ההיררכיה:
grid <- all the blocks of the kernel
└── block <- group of threads on the same SM
└── [warp] <- hardware implementation detail: 32 threads (not in the programming model!)
└── thread <- the atom: the code you write
שימו לב שכתבתי את ה-warp בסוגריים מרובעים ובהערה. זו לא קפריזה גרפית: היררכיית ה-threads של מודל התכנות היא thread -> block -> grid בלבד. ה-warp מוברח לתמונה כי החומרה מתעקשת עליו, אבל הוא לא שלב רשמי בהיררכיה שהמתכנת מבטא. נחזור לנקודה הזאת בהרחבה בהמשך; לעת עתה, זכרו את המבנה.
מבחינה חומרתית, ראינו ב-1.3 את המיפוי המדויק: "ליבת CUDA בודדת מריצה הוראות מ-thread בודד". כלומר ברמת ביצוע ההוראה, יש התאמה של אחד-לאחד בין ליבת CUDA ל-thread. אבל אל תטעו לחשוב שה-thread "שוכן" בליבה מסוימת: כפי שראינו ב-1.4, ה-thread חי בקובץ האוגרים של ה-SM, וה-warp scheduler מזרים אותו אל הליבות מחזור אחר מחזור. ה-thread הוא ישות תוכנתית; הליבה היא hardware pipeline שדרכו הוא זורם.
המצב הפרטי של thread - registers, stack ולמה אין syscalls¶
מה בדיוק "שייך" ל-thread? התשובה מפתיעה בצמצומה. בלשון הגלוסרי, כל thread "מחזיק את האוגרים שלו, אבל מעט מעבר לכך". המצב הפרטי של thread הוא, במהותו, רק ה-registers (אוגרים) שלו. אין ל-thread ערימת ניהול משלו, אין לו מרחב כתובות משלו, ואין לו הקשר של מערכת הפעלה. זהו הבדל תהומי מ-thread של POSIX ב-CPU.
יש ל-thread גם stack (מחסנית), אבל הוא יושב ב-global memory (זיכרון גלובלי), וקיומו נועד לשני צרכים בלבד:
- register spilling (שפיכת אוגרים) - כשה-thread צריך יותר אוגרים ממה שהוקצה לו, הcompiler "שופך" חלק מהמשתנים ל-stack. ראינו את ה-spill הזה בתרגול 1.4: הוא איטי, כי ה-stack חי בזיכרון הגלובלי האיטי.
- stack של קריאות פונקציה - כשה-kernel קורא לפונקציה שאינה inline, מסגרת הקריאה נשמרת ב-stack.
מכאן כלל אצבע מרכזי לכתיבת kernels מהירים: צמצמו את השימוש ב-stack. kernel מהיר שומר את כל המצב החם שלו באוגרים, נמנע מ-spilling (על ידי שמירת register pressure סביר, כפי שראינו ב-1.4), ונמנע מקריאות פונקציה עמוקות ומרקורסיה. כל בית שיורד ל-stack הוא בית שירד לזיכרון הגלובלי.
ההבדל החד ביותר מ-CPU: threads של CUDA אינם מבצעים syscalls (קריאות מערכת). thread של POSIX על CPU שייך ל-process (תהליך) שמערכת ההפעלה מתזמנת, יכול לפתוח קבצים, להקצות זיכרון, ולהיחסם על מנעולים של הליבה. thread של CUDA אינו כזה כלל: הוא אינו תהליך של מערכת הפעלה, הוא אינו מתוזמן על ידי גרעין ה-OS, ואין לו גישה ל-syscalls. הוא ישות הרבה יותר קלת-משקל - אין לו את הoverhead של הקשר OS, ולכן ה-GPU יכול להחזיק מאות אלפי threads חיים בו-זמנית (עד 2048 ל-SM ב-H100, כפי שראינו ב-1.4), מספר שבלתי-נתפס עבור threads של OS.
| מאפיין | thread של POSIX (CPU) | thread של CUDA (GPU) |
|---|---|---|
| מצב פרטי | אוגרים, stack מלא, מרחב כתובות, הקשר OS | בעיקר האוגרים; stack מינימלי ב-global |
| מי מתזמן | גרעין מערכת ההפעלה | ה-warp scheduler בחומרה (1.4) |
| syscalls | כן | לא |
| משקל | כבד (הקשר OS) | קל-משקל מאוד |
| כמה חיים בו-זמנית | עשרות עד אלפים למכונה | עד 2048 ל-SM, מאות אלפים ל-GPU |
| PC (program counter) פרטי | תמיד | אפשרי, אך lock-step הוא המסלול המהיר |
השורה האחרונה בטבלה מובילה לנקודה עדינה שראוי לחדד. ה-thread של CUDA יכול להחזיק PC (program counter - מונה תוכנית) פרטי משלו; כלומר, מבחינת מודל התכנות, threads מסוגלים לזרימת בקרה עצמאית. אבל לביצועים, kernels מהירים מריצים את כל 32 ה-threads של warp ב-lock-step (צעד-נעול), חולקים PC יחיד וזרם הוראות יחיד. זו נקודת המתח של SIMT: עצמאות אפשרית, אך יקרה; התכנסות (convergence) היא המסלול המהיר. וכאן בדיוק נכנס ה-warp.
ה-warp - יחידת ההרצה של החומרה¶
ה-warp הוא, בלשון הגלוסרי, "קבוצת threads שמתוזמנים יחד ומבוצעים במקביל". זו יחידת ההרצה הבסיסית על GPU של NVIDIA - הגרנולריות שבה החומרה באמת מתזמנת ומנפיקה הוראות. ראינו את זה מזווית החומרה ב-1.4: ה-warp scheduler אינו בוחר thread בודד למחזור, הוא בוחר warp שלם, ומנפיק את הוראתו לכל 32 ה-threads שלו בבת אחת.
גודל ה-warp הוא מספר קונקרטי שכדאי לחרוט: 32 threads ל-warp. הגלוסרי מנסח זאת בזהירות: "גודל ה-warp הוא טכנית קבוע תלוי-מכונה, אך בפועל (וגם בכל מקום אחר בגלוסרי הזה) הוא 32". כלומר, אף על פי שקיים המשתנה warpSize שהחומרה חושפת, ואף על פי ש-NVIDIA שומרת לעצמה את הזכות התיאורטית לשנותו, בכל דור GPU שקיים היום - כולל ה-H100 שהוא ה-GPU לדוגמה בקורס - הוא 32. בכל הקורס נניח 32, אך נכתוב קוד שקורא את warpSize במקום להטמיע את המספר, מטעמי היגיינה.
מכאן שני מספרים שנגזרים ישירות מ-1.4. ב-H100 יש ארבעה warp schedulers ל-SM, כל אחד מנפיק warp אחד למחזור:
warp = 32 threads
4 warp schedulers per SM -> 4 x 32 = 128 threads issued per SM per cycle
132 SMs in H100 -> 132 x 128 = 16,896 threads issued every clock cycle
resident cap: up to 64 warps = 2048 threads can reside simultaneously on an SM
ה-32 אינו מספר שרירותי. הוא תואם בדיוק את חלוקת ה-SM לארבע תת-מחיצות (sub-partitions) שראינו ב-1.4, שכל אחת מהן מחזיקה 32 ליבות FP32 - בדיוק רוחב warp אחד. כשמתזמן מנפיק הוראת FP32 של warp, 32 ה-threads שלו נופלים בול על 32 הליבות של תת-המחיצה, ורצים ב-lock-step. זה מוביל אותנו למודל ההרצה עצמו.
מודל ה-SIMT - Single Instruction, Multiple Threads¶
SIMT הוא ראשי התיבות של Single Instruction, Multiple Threads - הוראה יחידה, ריבוי threads. הגלוסרי מנסח: "כל ה-threads של warp מבצעים את אותה הוראה במקביל - המודל שנקרא Single-Instruction, Multiple Thread או SIMT". במילים אחרות, ה-warp הוא קבוצה של 32 threads שחולקים זרם הוראות אחד, וכל אחד מהם מפעיל אותה הוראה על נתונים שונים.
התיאור המדויק ביותר של המודל הזה מגיע דווקא מדף ה-PTX (Parallel Thread eXecution - הרצת שרשורים מקבילית, שפת הביניים שנעמיק בה בפרק 4) של הגלוסרי, ושווה לצטט אותו כמעט מילה במילה, כי הוא לוכד את כל הרעיון:
"במודל המכונה הזה יש יחידת הוראות אחת עבור מספר מעבדים. בעוד שכל מעבד מריץ thread אחד, ה-threads האלה חייבים לבצע את אותן הוראות - ומכאן הרצת שרשורים מקבילית, PTX. הם מתאמים זה עם זה דרך shared memory ומשיגים תוצאות שונות באמצעות האוגרים הפרטיים שלהם."
שלוש עובדות דחוסות במשפט הזה, וכולן מתחברות למה שכבר ראינו:
- יחידת הוראות אחת למספר מעבדים - יש שולף-ומפענח הוראות אחד (ה-warp scheduler מ-1.4), והוא מזין את כל 32 הpipelines של תת-המחיצה. זו החיסכון הגדול של ה-GPU: במקום 32 יחידות שליפה-ופענוח, יש אחת. הטרנזיסטורים שנחסכו הופנו לליבות חשבון נוספות.
- תיאום דרך shared memory - ה-threads מדברים זה עם זה דרך ה-shared memory (זיכרון משותף) של ה-block, כפי שנראה בפרק 3. שיתוף.
- בידול דרך registers פרטיים - זה מה שמאפשר ל-32 threads שמריצים את אותה הוראה להפיק תוצאות שונות: כל אחד מפעיל את ההוראה על הערכים באוגרים הפרטיים שלו. אותה הוראת
FMA, 32 קבוצות אוגרים, 32 תוצאות. בידול.
הנה המחשה מוחשית של warp שמבצע הוראת חיבור. שימו לב: הוראה אחת, 32 נתיבים (lanes), 32 תוצאות:
single instruction issued to the warp: c = a + b (on the private registers)
lane: 0 1 2 3 ... 31
a[lane]: 10 20 30 40 ... 310 (private registers)
b[lane]: 1 2 3 4 ... 32 (private registers)
--- --- --- --- ----
c[lane]: 11 22 33 44 ... 342 (32 different results)
^ ^ ^ ^ ^
|____ same instruction, same cycle, different data ____|
זהו ההבדל מ-SIMD (Single Instruction, Multiple Data) של CPU: ב-SIMD המתכנת מנהל וקטור רוחב-קבוע במפורש (למשל __m256), ואילו ב-SIMT המתכנת כותב קוד ל-thread סקלרי בודד, והחומרה היא זו שאורזת 32 threads ל-warp ומריצה אותם יחד. המודל התכנותי הוא סקלרי; המימוש הוא וקטורי. זו בדיוק הסיבה שאפשר לכתוב CUDA בלי להזכיר את המילה warp - כותבים thread אחד, וה-SIMT קורה מאחורי הקלעים.
ה-warp אינו חלק ממודל התכנות - an implementation detail¶
הגענו לתובנה החשובה ביותר של השיעור, וכדאי לומר אותה בפירוש כי מתחילים רבים מפספסים אותה. בלשון הגלוסרי: "warps אינם למעשה חלק מהיררכיית ה-threads של מודל התכנות של CUDA. במקום זאת, הם פרט מימוש של המימוש של המודל הזה על GPU של NVIDIA".
מה זה אומר בפועל? ההיררכיה הרשמית שמודל התכנות מבטיח היא thread -> block -> grid. אין בה warp. שום דבר בסמנטיקה של CUDA לא מבטיח לכם ש-threads מסוימים ירוצו יחד, או שהמספר 32 קדוש. קוד CUDA נכון - כזה שמפיק את התוצאה הנכונה - יעבוד גם אילו גודל ה-warp היה 16, 64, או 1. תוכלו לכתוב kernel שלם, לקמפל, ולהריץ, בלי לדעת בכלל שקיים דבר כזה warp, והתוצאה תהיה נכונה.
האנלוגיה שהגלוסרי מציע חדה ומדויקת: ה-warp הוא כמו שורת cache ב-CPU. שורת ה-cache אינה חלק ממודל התכנות של C. תוכנית C נכונה גם אם לא תדעו שקיים cache או שגודל השורה 64 בתים. אבל אם תתעלמו מ-cache lines, תכתבו קוד איטי - כי דילוג לא-מקומי בזיכרון קונס פספוסי cache יקרים. בדיוק כך עם warps: הם לא-רלוונטיים לנכונות, אבל חיוניים לביצועים. אפשר לכתוב CUDA נכון תוך התעלמות מ-warps; אי אפשר לכתוב CUDA מהיר תוך התעלמות מהם.
correctness performance
--------------------- ----------------------
thread essential essential
block essential essential
grid essential essential
warp not relevant <----------- fully essential
(cache line not relevant fully essential <- analogy from the CPU)
למה, אם כן, נקדיש שיעור שלם - ופרקים שלמים - ל-warp? כי כמעט כל אופטימיזציה שנלמד בהמשך היא, בבסיסה, "לחשוב במונחי warp". coalescing (איחוד גישות, פרק 3) הוא לסדר את גישות הזיכרון של 32 ה-threads של warp כך שיתמזגו לעסקה אחת. warp divergence (התפצלות ה-warp, שנפגוש מיד) היא כשה-threads של warp מתפצלים בזרימת הבקרה. bank conflicts (התנגשויות בנקים, פרק 3) הם כשה-threads של warp מתנגשים על בנקי ה-shared memory. הפרימיטיבים שנראה בסוף השיעור (__shfl_down_sync וחבריו) פועלים ברמת ה-warp. ה-warp הוא העדשה שדרכה מסתכלים על ביצועי GPU.
דירוג ה-warp ומספר ה-warps ל-block - warp rank¶
איך בעצם החומרה חותכת block ל-warps? הכלל פשוט ודטרמיניסטי: מחשבים את האינדקס הליניארי של ה-thread בתוך ה-block, ומחלקים ב-32. הגלוסרי נותן את הנוסחה במונחי ה-special registers של PTX; הנה היא, ולצידה ה-CUDA C++ המקביל:
// linear index of the thread within the block (row-major mapping of the 3 dimensions)
int linearIdx = threadIdx.x
+ threadIdx.y * blockDim.x
+ threadIdx.z * blockDim.x * blockDim.y;
int warpRank = linearIdx / warpSize; // which warp the thread belongs to (0,1,2,...)
int lane = linearIdx % warpSize; // position of the thread within the warp (0..31)
(ב-PTX אותו חישוב הוא linearIdx = %tid.x + %tid.y*%ntid.x + %tid.z*%ntid.x*%ntid.y; warpRank = linearIdx / 32;, כאשר %tid הם קואורדינטות ה-thread ו-%ntid הם ממדי ה-block. נראה את זה ב-PTX ממש בפרק 4.)
מספר ה-warps ל-block הוא עיגול כלפי מעלה של מספר ה-threads ל-block חלקי 32:
בואו נעבוד כמה מספרים קונקרטיים, כי כאן מסתתר מלכוד חשוב:
| גודל block | warps ל-block | threads פעילים ב-warp האחרון | lanes מבוזבזים |
|---|---|---|---|
| 32 | 1 | 32 | 0 |
| 256 | 8 | 32 | 0 |
| 1024 (המקסימום) | 32 | 32 | 0 |
| 96 | 3 | 32 | 0 |
| 100 | 4 | 4 | 28 |
| 33 | 2 | 1 | 31 |
שתי השורות האחרונות הן המלכוד: block של 100 threads אינו נחתך ל"3.125 warps" - הוא נחתך ל-4 warps, כשה-warp הרביעי מכיל 4 threads פעילים ו-28 נתיבים מבוזבזים. אותם 28 נתיבים עדיין נלקחים בחשבון בתזמון ובמשאבים, אבל אינם עושים עבודה מועילה. לכן הכלל המעשי: בחרו גודל block שהוא כפולה של 32 (32, 64, 128, 256, ...). block של 33 threads הוא בזבוז חמור - 2 warps שלמים כדי להריץ 33 threads, כש-31 נתיבים ריקים.
דוגמה דו-ממדית שממחישה את המיפוי הליניארי. נניח block בגודל (16, 16) = 256 threads. ה-thread בקואורדינטות (threadIdx.x=5, threadIdx.y=3):
linearIdx = 5 + 3*16 = 53
warpRank = 53 / 32 = 1 (the second warp, warp number 1)
lane = 53 % 32 = 21 (lane 21 within that warp)
שימו לב לתוצאה החשובה של המיפוי row-major: ב-block דו-ממדי, ה-warp הראשון (rank 0) מכסה את threads (0..15, 0) ואת (0..15, 1) - כלומר שתי שורות שלמות של 16, כי 16 + 16 = 32. זו עובדה קריטית ל-coalescing: אם השורה הסמוכה בזיכרון היא ציר ה-x, אז 32 ה-threads של warp ניגשים ל-32 אלמנטים רצופים - בדיוק מה שרוצים. מיפוי ה-warps לקואורדינטות אינו שרירותי, והוא ישפיע ישירות על ביצועים.
לראות ולנצל את ה-warp - warp-level primitives¶
עד כה אמרנו ש-warps הם "מאחורי הקלעים". אבל CUDA נותנת פרימיטיבים מפורשים שמאפשרים ל-threads של warp לתקשר ישירות, בלי לעבור דרך shared memory - וזו אחת מיכולות הביצועים החזקות של המודל. שלושה שנפגוש בתרגול:
__activemask()- מחזיר מסכה של 32 ביט שבה כל ביט דולק מציין נתיב פעיל ב-warp הנוכחי. זו הדרך לראות "מי בבית" ב-warp שלי ברגע נתון.__ballot_sync(mask, predicate)- כל thread תורם ביט (הערך שלpredicateאצלו), והתוצאה היא מסכת 32 ביט מכל ה-warp. הצבעה קולקטיבית של ה-warp במחיר הוראה אחת.__shfl_down_sync(mask, val, delta)- "שאיבה" (shuffle): ה-thread מקבל את הערךvalמה-thread שנמצאdeltaנתיבים מעליו, ישירות מהאוגר שלו, בלי shared memory. זהו הבסיס ל-reduction (צמצום) ברמת warp.
הנה ה-worked example המרכזי: reduction ברמת warp, שסוכם את 32 הערכים של warp לערך בודד בעזרת חמש הוראות shuffle בלבד (log2(32) = 5), בלי shared memory ובלי __syncthreads():
// sums the value of all 32 threads in the warp; the result ends up in lane 0
__inline__ __device__ float warpReduceSum(float val) {
// 16 -> 8 -> 4 -> 2 -> 1 : five steps, each one folds the warp in half
for (int offset = warpSize / 2; offset > 0; offset >>= 1) {
val += __shfl_down_sync(0xffffffff, val, offset);
}
return val; // lane 0 holds the full sum of the warp
}
המסכה 0xffffffff היא 32 ביט דלוקים - "כל הנתיבים משתתפים". חובה לספק אותה מפורשות: מאז דור Volta (ראו הערה על independent thread scheduling בהמשך), הפרימיטיבים ה"מסונכרנים" (סיומת _sync) דורשים מסכה שמצהירה אילו נתיבים חייבים להיפגש בהוראה - וזו הגנה מפני באגים עדינים של warp מתפצל. הנה זרימת הקיפול:
offset=16: lane i += lane i+16 (32 values -> 16 partial sums)
offset=8: lane i += lane i+8 (16 -> 8)
offset=4: lane i += lane i+4 (8 -> 4)
offset=2: lane i += lane i+2 (4 -> 2)
offset=1: lane i += lane i+1 (2 -> 1) <- lane 0 = the full sum
זה מהיר בהרבה מ-reduction דרך shared memory, כי הנתונים אף פעם לא עוזבים את האוגרים. זו דוגמה מובהקת ל"לחשוב במונחי warp": התוצאה של סכימת 32 מספרים הייתה נכונה גם עם לולאה תמימה, אבל המהירות תלויה לגמרי בניצול מבנה ה-warp.
התפצלות ה-warp - warp divergence (הצצה). מה קורה כשלא כל ה-threads של warp רוצים ללכת לאותו מקום? נניח if (threadIdx.x % 2 == 0) { ... } else { ... }. חצי מה-warp רוצה את ענף ה-if וחצי את ה-else, אבל ל-warp יש PC אחד. הפתרון של החומרה: להריץ את שני הענפים בזה אחר זה, כשבכל ענף חלק מהנתיבים "מכובים" (predicated off). בלשון הגלוסרי: "כשה-threads ב-warp מתפצלים זה מזה כדי לבצע הוראות שונות, המצב הידוע גם כ-warp divergence, הביצועים נופלים בדרך כלל בחדות". במקרה הגרוע, warp שמתפצל ל-32 מסלולים שונים רץ פי 32 לאט יותר. זו אחת מסכנות הביצועים המרכזיות ב-GPU, ונקדיש לה טיפול מלא בשיעור 8.3 - כולל השינוי החשוב של Volta (independent thread scheduling), שבו כל thread כן מקבל PC פרטי משלו, מה שמשנה את הדינמיקה של ההתכנסות. לעת עתה, זכרו את הכלל: נסו לגרום לכל 32 ה-threads של warp ללכת לאותו מסלול.
מעבר ל-API אחרים ואטימולוגיה¶
הרעיון של "קבוצת threads שרצה יחד ב-lock-step" אינו ייחודי ל-CUDA; כל ה-API של תכנות GPU הגיעו לאותה מסקנה חומרתית, ורק נתנו לה שמות שונים. שווה להכיר את המילון, כי ספרות ותיעוד חוצי-פלטפורמות ישתמשו בשמות האלה:
| API / פלטפורמה | השם ליחידה של threads רצים-יחד |
|---|---|
| NVIDIA CUDA | warp (32 threads) |
| WebGPU | subgroups |
| DirectX (HLSL) | waves |
| Apple Metal | simdgroups |
כולם מתארים את אותו הדבר: הגרנולריות שבה החומרה מנפיקה הוראה אחת למספר threads. אם תיתקלו ב"wave intrinsics" ב-HLSL או ב"subgroup operations" ב-WebGPU, אלו המקבילות המדויקות ל-warp-level primitives שראינו כאן.
ולסיום, מהיכן השם המוזר "warp"? הוא אינו קשור ל"warp speed" של מסע בין כוכבים. השם בא מעולם האריגה. ב-Lindholm et al., 2008 - אותו מאמר לבן שתיעד את החזון המקורי של CUDA - האריגה מתוארת כ"טכנולוגיית ה-threads המקבילה הראשונה". בנול אריגה, ה-warp הוא אוסף הthreads (threads!) המקבילים המתוחים לאורך, שדרכם עובר thread השתי. המטפורה מושלמת: אוסף של threads מקבילים המוחזקים במתח ומעובדים יחד. גם המילה האנגלית thread עצמה, "thread", מגיעה מאותו עולם סמנטי. השמות בתחום נבחרו בקפידה, וכדאי לזכור אותם - הם עוזרים לחשוב.
סיכום¶
- ה-thread הוא האטום של היררכיית ה-threads של מודל התכנות של CUDA (thread -> block -> grid); הקוד שאתם כותבים ב-
__global__הוא הקוד של thread בודד, והחומרה מריצה ממנו עותקים רבים. - המצב הפרטי של thread הוא בעיקרו ה-registers שלו בלבד; יש לו גם stack ב-global memory לצורך register spilling ולקריאות פונקציה, ו-kernels מהירים ממזערים את השימוש ב-stack האיטי.
- ה-threads של CUDA אינם מבצעים syscalls ואינם תהליכי OS; הם קלי-משקל בהרבה מ-threads של POSIX, ולכן GPU יכול להחזיק עד 2048 threads תושבים ל-SM ומאות אלפים ל-GPU.
- ה-warp הוא קבוצת 32 threads שהחומרה מתזמנת ומנפיקה יחד - יחידת ההרצה הבסיסית של NVIDIA; גודל ה-warp הוא טכנית תלוי-מכונה אך בפועל 32 בכל דור, כולל H100.
- מודל ה-SIMT (Single Instruction, Multiple Threads) מדויק במודל המכונה של PTX: יחידת הוראות אחת למספר מעבדים, ה-threads מבצעים את אותה הוראה, מתאמים דרך shared memory, ומשיגים תוצאות שונות דרך האוגרים הפרטיים.
- ה-warp אינו חלק ממודל התכנות הרשמי - הוא פרט מימוש של NVIDIA, אנלוגי לשורת cache ב-CPU: לא-רלוונטי לנכונות, חיוני לחלוטין לביצועים; אפשר לכתוב CUDA נכון בלי warps, אך לא CUDA מהיר.
- ה-warp rank של thread הוא
linearIdx / 32ומספר ה-warps ל-block הואceil(threads/32); block שאינו כפולה של 32 מבזבז נתיבים ב-warp האחרון, ולכן בוחרים גודל block שהוא כפולה של 32. - פרימיטיבים ברמת warp (
__activemask(),__ballot_sync(),__shfl_down_sync()) מאפשרים ל-threads של warp לתקשר ישירות דרך האוגרים; reduction ברמת warp מסכם 32 ערכים בחמישה shuffles בלבד, בלי shared memory. - ה-warp divergence היא כשה-threads של warp מתפצלים בזרימת הבקרה, וה-warp מריץ את הענפים בזה אחר זה - נפילת ביצועים חדה שנטפל בה במלואה ב-8.3; הכלל: השאירו את 32 ה-threads של warp על מסלול אחד.
- אותה יחידה נקראת subgroups ב-WebGPU, waves ב-DirectX ו-simdgroups ב-Metal; השם "warp" בא מאריגה (Lindholm 2008) - אוסף threads מקבילים מתוחים במקביל.