3.4 זיכרון קבוע וטקסטורה תרגול

תרגול - זיכרון קבוע וטקסטורה

בתרגול הזה תיישמו בידיים את שלושת מרחבי הזיכרון המיוחדים שראינו בהרצאה, ותמדדו את ההשפעה שלהם. תתחילו מסטנסיל 1D שמקדמיו בזיכרון קבוע ותשוו אותו לגרסה שקוראת אותם מ-global memory, תנתבו קלט לקריאה-בלבד עם const __restrict__ ותסתכלו ב-SASS כדי לוודא שקיבלתם LDG.E.CI, תדגמו תמונה עם אובייקט טקסטורה ואינטרפולציה בי-לינארית בחומרה, ולבסוף תראו במו עיניכם מתי זיכרון קבוע מזיק - כשהכתובות מתפצלות. עבדו לפי הסדר. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס, אבל כל תרגיל כאן ירוץ על כל GPU של NVIDIA מדור Kepler ומעלה.

הכנה

צרו קובץ stencil.cu. השתמשו במקרו CUDA_CHECK מההרצאה 0.3 בראש הקובץ (הוא יחזור בכל תרגיל). הקומפילציה לאורך התרגול (החליפו את sm_90a בארכיטקטורה שלכם - למשל sm_75 ל-T4):

nvcc -O2 -arch=sm_90a -o stencil stencil.cu

לברור ה-compute capability של הכרטיס:

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

לכל המדידות בתרגול השתמשו באירועי CUDA (cudaEvent) כפי שראינו בשיעור 0.3, והריצו הרצת "חימום" אחת לפני המדידה כדי לא לכלול overhead של compilation-בזמן-ריצה חד-פעמית.

תרגיל 1 - מקדמי סטנסיל בזיכרון קבוע מול global

1. ממשו סטנסיל 1D עם RADIUS = 4 (כלומר DIAM = 9 מקדמים). כתבו שתי גרסאות של ה-kernel: stencil_const שקורא את המקדמים מ-__constant__ float c_coef[DIAM], ו-stencil_global שמקבל את המקדמים כמצביע const float* coef ב-global memory.
2. אתחלו את המקדמים למסנן ממוצע פשוט (1.0f / DIAM בכל תא). טענו אותם לזיכרון הקבוע עם cudaMemcpyToSymbol, ולגרסת ה-global עם cudaMemcpy רגיל.
3. הריצו את שתי הגרסאות על מערך של n = 1 << 26 (כ-67 מיליון איברים), אמתו שהפלט זהה בין השתיים, ומדדו את זמן ה-kernel של כל אחת עם cudaEvent.
4. חשבו את יחס המהירות והסבירו אותו: מדוע הגרסה עם זיכרון קבוע מהירה (או לפחות לא איטית) יותר?

רמז: בלולאה הפנימית האינדקס למקדם (k + RADIUS) אחיד על פני כל ה-warp - כל 32 הthreads קוראים את אותו c_coef[...] באותו רגע. זו תבנית ה-broadcast: קריאה אחת משרתת את כל ה-warp. בגרסת ה-global, אותה קריאה עוברת דרך נתיב ה-L1 הרגיל בלי מנגנון ה-broadcast הייעודי. ההפרש עשוי להיות קטן כאן כי המקדמים חמים ב-L1 בשתי הגרסאות - החשיבות של constant תתחדד בתרגיל 4.

תרגיל 2 - קלט לקריאה-בלבד ובדיקת SASS

1. הוסיפו ל-stencil_const גרסה stencil_ro שבה מצביע הקלט מסומן const float* __restrict__ in (וגם מצביע הפלט float* __restrict__ out).
2. קמפלו וחלצו את ה-SASS של שני ה-kernels:

nvcc -O2 -arch=sm_90a -cubin -o stencil.cubin stencil.cu
cuobjdump -sass stencil.cubin | grep -E "LDG|LDC"

1. חפשו את ההבדל: בגרסה עם const __restrict__ טעינות הקלט אמורות להופיע כ-LDG.E.CI (הנתיב לקריאה-בלבד), ובגרסה בלי - כ-LDG.E רגיל.
2. שימו לב גם היכן מופיעים המקדמים מזיכרון הקבוע: הם לרוב לא LDG בכלל, אלא LDC או אופרנד מסוג c[0x3][...] (בנק קבוע).

רמז: אם שתי הגרסאות מפיקות LDG.E.CI גם בלי __restrict__, ייתכן שהcompiler הסיק לבד שהקלט לקריאה-בלבד (זה קורה כשאין כתיבה נראית לאותו מצביע). כדי לראות את ההבדל בבירור, נסו גרסה שבה מצביע הקלט אינו const, או השתמשו ב-__ldg(&in[i+k]) במפורש ובדקו שה-LDG.E.CI צץ. הסיומת .CI = Cache in Invalidate, כלומר cache הקריאה-בלבד הלא-קוהרנטי.

תרגיל 3 - דגימת תמונה עם אינטרפולציה בי-לינארית בחומרה

1. צרו קובץ חדש resample.cu. בנו תמונת קלט קטנה סינתטית בגודל W x H (למשל 4x4) עם ערכים ידועים - למשל gradient לינארי img[y*W + x] = (float)(x + y).
2. הקימו cudaTextureObject_t מעל cudaArray, עם filterMode = cudaFilterModeLinear, addressMode = cudaAddressModeClamp, ו-normalizedCoords = 0, בדיוק כפי שראינו בהרצאה.
3. כתבו kernel שדוגם את התמונה בנקודות ביניים - למשל בקואורדינטות (1.0f, 1.0f), (1.5f, 1.5f), (2.5f, 1.5f). זכרו שמרכז טקסל (px, py) נמצא בקואורדינטה (px + 0.5f, py + 0.5f).
4. הדפיסו את הערכים שהחזירה הדגימה ואמתו ידנית: דגימה בדיוק במרכז טקסל צריכה להחזיר את ערך הטקסל, ודגימה באמצע בין שני טקסלים צריכה להחזיר את ממוצעם.

רמז: בקואורדינטה (1.0f, 1.0f) (לא במרכז טקסל) אתם באמצע בין ארבעה טקסלים ותקבלו את הממוצע שלהם. עם img[y*W+x] = x + y, הטקסל (1,1) שווה 2, (0,0) שווה 0, וכו'; הדגימה ב-(1.0,1.0) (פינת ארבעת הטקסלים (0,0),(1,0),(0,1),(1,1)) תיתן (0+1+1+2)/4 = 1.0. אל תשכחו לשחרר עם cudaDestroyTextureObject ו-cudaFreeArray. אם קיבלתם אפסים, ודאו שאתחלתם את resDesc ו-texDesc ל-{} (אתחול לאפס) לפני מילוי השדות.

תרגיל 4 - מתי זיכרון קבוע מזיק - broadcast מול כתובות מתפצלות

1. בנו טבלה קטנה בזיכרון קבוע: __constant__ float c_table[256];, מאותחלת לערכים כלשהם.
2. כתבו שני kernels שכל אחד קורא איבר אחד מהטבלה לכל thread וכותב אותו לפלט:
3. read_broadcast: כל הthreads קוראים את אותה כתובת, c_table[0].
4. read_divergent: כל thread קורא כתובת אחרת, c_table[threadIdx.x & 255].
5. הריצו את שניהם על n = 1 << 26, מדדו את זמן ה-kernel של כל אחד עם cudaEvent, וחשבו את היחס ביניהם.
6. הסבירו את התוצאה במונחי מספר הטרנזקציות לכל warp. מה היה קורה אילו c_table היה ב-global memory מאוחד במקום בזיכרון קבוע?

רמז: ב-read_broadcast כל warp מבקש כתובת אחת -> טרנזקציה אחת; ב-read_divergent כל warp מבקש עד 32 כתובות שונות -> עד 32 טרנזקציות מסודרות בטור. היחס יהיה משמעותי (עד פי כמה). לו הטבלה הייתה ב-global memory, ה-read_divergent דווקא היה מהיר יותר, כי 32 הכתובות threadIdx.x & 255 של warp רציפות ומתאחדות (coalesce) לטרנזקציה אחת - זו בדיוק ההמחשה לכך שזיכרון קבוע ממוטב ל-broadcast ולא לגישה מתפצלת.

תרגיל 5 (בונוס) - סטנסיל 2D: טקסטורה מול global

1. ממשו את הסטנסיל הדו-ממדי הממוצע 5-נקודתי (jacobi5pt) בשתי גרסאות: אחת שקוראת מ-cudaTextureObject_t (עם cudaAddressModeClamp שמטפל בגבולות), ואחת שקוראת מ-global memory עם בדיקות גבול ידניות (if (x > 0) וכו').
2. הריצו את שתיהן על תמונה בגודל 4096 x 4096, אמתו שהפלט זהה (עד הבדלי גבול), ומדדו את זמן ה-kernel של כל אחת.
3. חשבו את רוחב הפס האפקטיבי של כל גרסה (כל פיקסל קורא 5 ערכים וכותב 1, כלומר 6 * 4 = 24 בתים) והשוו לרוחב הפס התיאורטי של הכרטיס.
4. הסיקו: מתי הייתם בוחרים בכל אחת מהגרסאות? חשבו גם על גרסת shared-memory tiling מ-3.3 - איך היא הייתה משתלבת?

רמז: הcache הטקסטורה עם הלוקליות הדו-ממדית שלו לרוב מנצח את global ה"נאיבי" כשהגישה האנכית (yc - 1, yc + 1) גורמת ל-cache misses ב-global. עם זאת, גרסת shared-memory tiling (טעינת tile משותף פעם אחת ל-block) בדרך כלל מנצחת את שתיהן בחומרה מודרנית - ה-texture מנצחת בעיקר בקלות הכתיבה ובטיפול-הגבולות-החינמי. את המסגרת המלאה של רוחב פס וניתוח roofline נעמיק בפרק 7.