3.4 זיכרון קבוע וטקסטורה הרצאה

בשיעור 3.1 מיפינו את היררכיית הזיכרון הלוגית של CUDA - registers פרטיים ל-thread, shared memory (זיכרון משותף) ברמת ה-block, ו-global memory (זיכרון גלובלי) שכל ה-grid רואה - והדגשנו שהיא היררכיה מנוהלת-מתכנת, לא cache שקוף כמו ב-CPU. בשיעור 3.3 ראינו איך shared memory הופך kernel memory-bound למהיר בעזרת tiling - טעינת tile נתונים פעם אחת מ-global memory ושימוש-חוזר בו הרבה. בשיעור הזה נסגור את התמונה עם שלושה מרחבי זיכרון מיוחדים שכולם משרתים קריאה בלבד: זיכרון קבוע (constant memory), cache הקריאה-בלבד (read-only data cache), וזיכרון טקסטורה (texture memory). כל אחד מהם הוא נתיב חומרה ייעודי שממוטב לתבנית גישה מסוימת - שידור (broadcast) של אותו ערך לכל ה-warp, שימוש-חוזר של נתונים קבועים, או מרחביות דו-ממדית עם אינטרפולציה בחומרה. שלושתם אינם מחליפים את global ו-shared, אלא נותנים לנו כלים חדים יותר כשתבנית הקריאה מתאימה. עם הכלים האלה ביד נהיה מוכנים לפרויקט 3.5, שם נצטרך לבחור נכון בין כל המרחבים.

למה נתיב-קריאה נפרד - the read-side memory spaces

ב-3.1 ראינו שגישה ל-global memory יקרה - מאות מחזורי שעון של latency (השהיה) לכל טעינה שלא נמצאת בcache. ה-tiling ב-shared memory פותר את זה כשיש שימוש-חוזר בתוך block, אבל יש תבניות קריאה שבהן shared memory מסורבל או מיותר, ובכל זאת global memory "טהור" מבזבז. שלוש דוגמאות קלאסיות:

• מקדם שכל הthreads קוראים בו-זמנית. במסנן קונבולוציה, כל 32 הthreads של warp מכפילים את שכנם במקדם w[k] - אותו מקדם בדיוק. לקרוא אותו מ-global memory פירושו 32 threads שמושכים את אותה כתובת. זה בזבוז: היינו רוצים חומרה שמזהה "כולם רוצים את אותו דבר" ומשרתת את כולם בפעולה אחת. זה בדיוק constant memory.
• מערך קלט לקריאה-בלבד עם שימוש-חוזר. אם אנחנו יודעים בוודאות שמצביע מסוים נקרא ולעולם לא נכתב בתוך ה-kernel, אפשר לנתב אותו דרך cache ייעודי לקריאה-בלבד, שלא צריך לשלם על מנגנוני קוהרנטיות של כתיבה. זה cache הקריאה-בלבד (__ldg).
• גישה מרחבית דו-ממדית עם אינטרפולציה. בעיבוד תמונה, בסטנסילים ובדגימה מחדש (resampling), threads שכנים ניגשים לפיקסלים שכנים, ולעיתים לקואורדינטות שאינן שלמות ודורשות אינטרפולציה. חומרת ה-texture בנויה בדיוק לזה: הcache עם לוקליות דו-ממדית, טיפול בגבולות, ואינטרפולציה בי-לינארית בחומרה.

המשותף לשלושתם: הם קריאה בלבד מנקודת המבט של ה-kernel (constant ו-texture חסומים לכתיבה מ-device), והם נשענים על caches ייעודיים בתוך ה-SM (מעבד רב-זרמי) שאינם על נתיב הכתיבה הרגיל. נעבור עליהם אחד-אחד.

זיכרון קבוע - constant memory

זיכרון קבוע הוא מרחב בגודל 64 KiB לכל מודול/device, שמצהירים עליו עם המזהה __constant__ בהיקף קובץ (file scope), ושה-kernel יכול רק לקרוא ממנו. פיזית הנתונים יושבים ב-GPU RAM כמו global memory, אבל הם ממוטמנים בcache קבוע ייעודי בכל SM (per-SM constant cache) - cache קטן, בסדר גודל של קילובייטים בודדים לכל SM. הייחוד שלו הוא בתבנית ה-broadcast (שידור): כשכל 32 הthreads של warp קוראים את אותה כתובת, החומרה משרתת את כולם בפעולה אחת, במחיר של קריאת register כמעט. זו הסיבה שהוא מושלם למקדמים, לפרמטרים ולקבועים שכל הthreads חולקים.

מצהירים ומאתחלים כך. שימו לב: המשתנה מוצהר גלובלית מחוץ לכל פונקציה, וטוענים אליו מהhost עם cudaMemcpyToSymbol - לא עם cudaMemcpy רגיל:

#define RADIUS 4
#define DIAM   (2 * RADIUS + 1)          // 9 coefficients

__constant__ float c_coef[DIAM];         // lives in constant memory, read-only for the kernel

__global__ void stencil1d(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= RADIUS && i < n - RADIUS) {
        float acc = 0.0f;
        #pragma unroll
        for (int k = -RADIUS; k <= RADIUS; ++k) {
            acc += c_coef[k + RADIUS] * in[i + k];   // c_coef[k+RADIUS] is identical for the whole warp
        }
        out[i] = acc;
    }
}

מהhost, לפני הlaunch:

float h_coef[DIAM] = { /* ... filter coefficients ... */ };
CUDA_CHECK(cudaMemcpyToSymbol(c_coef, h_coef, sizeof(h_coef)));

שימו לב לכמה דקויות חשובות:

• מעבירים את שם הסמל, לא מצביע. הארגומנט הראשון של cudaMemcpyToSymbol הוא c_coef עצמו (הסמל), לא &c_coef. ה-runtime מתרגם את הסמל לכתובת בזיכרון הקבוע.
• הגודל 64 KiB הוא גבול קשיח. אם תכריזו על __constant__ שחורג מ-64 KiB, הקומפילציה תיכשל. הזיכרון הקבוע נועד לכמויות קטנות - מקדמים, טבלאות פרמטרים - לא למערכי נתונים גדולים.
• __constant__ אינו כמו const. זה מרחב זיכרון פיזי עם נתיב cache משלו, לא סתם מבטיח-אי-שינוי של C++.
• פרמטרים של kernel גם הם בזיכרון קבוע. הארגומנטים שמעבירים ב-<<<>>>(...) (כמו n והמצביעים) יושבים בבנק קבוע נפרד ומגיעים דרך אותו מנגנון broadcast - זו הסיבה שקריאת פרמטר ב-kernel זולה כל כך.

בקוד למעלה, בכל איטרציה של הלולאה, k אחיד (uniform) על פני ה-warp - כל 32 הthreads באותו k באותו רגע - ולכן c_coef[k + RADIUS] הוא אותה כתובת לכל ה-warp. זו תבנית ה-broadcast המושלמת.

תבנית השידור מול כתובות מתפצלות - broadcast vs divergent addresses

הכוח של constant memory תלוי לחלוטין בתבנית הגישה, ולכן חובה להבין מתי הוא עוזר ומתי הוא מזיק. החוקיות פשוטה: הזמן לקרוא מזיכרון קבוע פרופורציוני למספר הכתובות השונות ש-warp מבקש בו-זמנית.

   reading from constant memory in one warp (32 threads)
   ┌───────────────────────────────────────────────────────────┐
   │  all 32 threads read address A          →  1 operation       │  perfect broadcast
   │  16 threads read A, 16 read B           →  2 serialized ops  │
   │  each thread reads a different address (32 different) →  32 serialized ops │  the worst case
   └───────────────────────────────────────────────────────────┘

בואו נתרגם למספרים. נניח warp שקורא ערך יחיד מזיכרון קבוע:

• broadcast מלא: כתובת אחת שונה -> פעולה אחת. תפוקה אפקטיבית: 32 threads משורתים במחיר של קריאה בודדת.
• התפצלות מלאה: 32 כתובות שונות -> 32 פעולות מסודרות בטור (serialized). התוצאה פי 32 איטית יותר מה-broadcast, ובמקרה כזה global memory מאוחד (coalesced) - שקורא 32 כתובות רציפות בטרנזקציה אחת - יהיה מהיר בהרבה.

מכאן הכלל: זיכרון קבוע מנצח כשהאינדקס לתוכו אחיד על פני ה-warp, ומפסיד כשהאינדקס תלוי ב-threadIdx. בסטנסיל למעלה האינדקס k + RADIUS אחיד, ולכן זה ניצחון. אבל שקלו את ה-kernel ההפוך:

// anti-pattern: each thread reads a different address in constant memory
__global__ void bad_lookup(const int* idx, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        out[i] = c_table[idx[i] % DIAM];   // idx[i] differs between threads -> divergent addresses
    }
}

כאן idx[i] שונה בין threads סמוכים, ולכן הthreads ב-warp מבקשים כתובות שונות בזיכרון הקבוע, והחומרה מסדרת אותן בטור. זהו בדיוק המצב שבו constant memory מזיק - טבלת חיפוש שנקראת באינדקסים מתפצלים שייכת ל-global memory (או ל-texture), לא לזיכרון הקבוע. בתרגול תמדדו את שני המקרים ותראו את יחס ה-32 בעיניכם.

הcache הקריאה-בלבד - the read-only data cache

לא כל נתון קבוע ראוי או נכנס ל-64 KiB. לרוב יש לנו מערך גדול ב-global memory שאנחנו יודעים שהוא נקרא בלבד בתוך ה-kernel. במקרה כזה אפשר לנתב את הטעינות דרך cache הקריאה-בלבד (read-only data cache, המכונה גם texture cache) - נתיב cache שאינו צריך לתחזק קוהרנטיות של כתיבה, וממוטב לנתונים שרק נקראים. יש שתי דרכים לבקש אותו:

// way 1: explicit intrinsic
out[i] = __ldg(&in[i]);

// way 2: promise the compiler with const __restrict__ (recommended - more readable)
__global__ void stencil_ro(const float* __restrict__ in,
                           float* __restrict__ out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= RADIUS && i < n - RADIUS) {
        float acc = 0.0f;
        #pragma unroll
        for (int k = -RADIUS; k <= RADIUS; ++k)
            acc += c_coef[k + RADIUS] * in[i + k];   // the compiler may route this through the read-only cache
        out[i] = acc;
    }
}

הצירוף const __restrict__ אומר לcompiler שני דברים: const - הנתון לא נכתב דרך המצביע הזה; __restrict__ - אף מצביע אחר לא מצביע לאותו זיכרון (no aliasing). בהינתן שני אלה, ה-ptxas (הcompiler מ-PTX ל-SASS) חופשי להנפיק טעינה דרך cache הקריאה-בלבד. ב-SASS זה מופיע כ-LDG.E.CI (הסיומת .CI = Cache in Invalidate, כלומר הנתיב הלא-קוהרנטי לקריאה-בלבד), בניגוד ל-LDG.E הרגיל:

   regular read from global:        LDG.E   R4, [R2]        ; regular L1 path
   read through the read-only cache: LDG.E.CI R4, [R2]        ; the non-coherent path

מתי זה עוזר, ומתי לא:

• עוזר כשהנתון נקרא-בלבד ויש שימוש-חוזר או גישה בלתי-מאוחדת (non-coalesced) שנהנית מהcache. הוא גם מוריד לחץ מנתיב ה-L1 הרגיל.
• פחות קריטי בחומרה מודרנית. מ-Volta והלאה (וכך ב-Hopper, ה-H100 שלנו) ה-L1 וה-texture cache מאוחדים לאותה SRAM, ולכן ההפרדה הפיזית שהייתה משמעותית ב-Kepler פחות דרמטית. עדיין, const __restrict__ נותן לcompiler מידע יקר (אי-aliasing) שמאפשר לו לתזמן טעינות אגרסיבי יותר - אז זו הרגל טוב בכל מקרה.
• לא רלוונטי לנתון שנכתב. אם המצביע נכתב אי-שם ב-kernel, אסור לסמן אותו const, וcache הקריאה-בלבד עלול להחזיר נתון מיושן.

הדרך הפרקטית לוודא שקיבלתם את הנתיב שרציתם היא להסתכל ב-SASS - בדיוק מה שנעשה בתרגול עם cuobjdump -sass ו-grep LDG.

זיכרון טקסטורה - texture memory

זיכרון הטקסטורה הוא הנתיב המתוחכם ביותר בשלושה. מקורו בגרפיקה - מיפוי טקסטורות על משטחים - ולכן החומרה שלו בנויה סביב לוקליות מרחבית דו-ממדית ודגימה (sampling) של הנתון, לא רק קריאה שלו. הוא ניגש דרך אובייקט טקסטורה, cudaTextureObject_t (ה-API המודרני, מ-CUDA 5.0 / Kepler והלאה, שהחליף את ה-texture references הישנים והמסורבלים). הוא נותן לנו בחומרה, בחינם, ארבע יכולות שקשה לשחזר ידנית:

• הcache עם לוקליות דו-ממדית. בניגוד לcache קווי שממטמן שורות רציפות, cache הטקסטורה ממטמן tiles דו-ממדיים - אז threads שכנים שקוראים פיקסלים שכנים (למעלה/למטה/לצדדים) נהנים ממנו, גם כשהגישה אינה מאוחדת בזיכרון הקווי.
• קואורדינטות מנורמלות - normalized coordinates. אפשר לפנות לתמונה בקואורדינטות בטווח [0, 1) במקום באינדקסים שלמים, כך שאותו kernel עובד לכל רזולוציה.
• מצבי גבול - addressing modes. מה קורה כשפונים מחוץ לתמונה: cudaAddressModeClamp (הידוק לקצה), cudaAddressModeWrap (גלישה מחזורית), cudaAddressModeMirror (שיקוף), cudaAddressModeBorder (אפס מחוץ לגבול). זה מבטל את כל בדיקות הגבולות הידניות בסטנסיל.
• אינטרפולציה בי-לינארית בחומרה - hardware linear interpolation. עם cudaFilterModeLinear, כשפונים לקואורדינטה לא-שלמה, החומרה מחזירה שקלול בי-לינארי של ארבעת הטקסלים השכנים - בפעולת חומרה אחת. משקלי האינטרפולציה מחושבים בדיוק שבור של 8 ביטים (fixed-point), כך שזו אינטרפולציה מהירה אך לא מדויקת-להפליא - מצוין לגרפיקה ולעיבוד תמונה, פחות לחישוב מדעי שדורש דיוק מלא.

הקמת אובייקט טקסטורה דורשת שרשרת של מבני תיאור. הזרימה: מקצים cudaArray (מבנה זיכרון מיוחד עם פריסה דו-ממדית ממוטבת ל-texture), מעתיקים אליו את התמונה, ואז בונים resource descriptor (מה הנתון) ו-texture descriptor (איך לדגום אותו):

int W = ..., H = ...;                    // image dimensions
// 1. allocate cudaArray on the device and copy the image into it
cudaChannelFormatDesc ch = cudaCreateChannelDesc<float>();
cudaArray_t cuArr;
CUDA_CHECK(cudaMallocArray(&cuArr, &ch, W, H));
CUDA_CHECK(cudaMemcpy2DToArray(cuArr, 0, 0, h_img,
                               W * sizeof(float),      // pitch of the source
                               W * sizeof(float), H,   // width in bytes, height
                               cudaMemcpyHostToDevice));

// 2. resource descriptor: what the data is
cudaResourceDesc resDesc = {};
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cuArr;

// 3. texture descriptor: how to sample it
cudaTextureDesc texDesc = {};
texDesc.addressMode[0] = cudaAddressModeClamp;   // x axis
texDesc.addressMode[1] = cudaAddressModeClamp;   // y axis
texDesc.filterMode     = cudaFilterModeLinear;   // hardware bilinear interpolation
texDesc.readMode       = cudaReadModeElementType;
texDesc.normalizedCoords = 0;                    // coordinates in pixels, not in [0,1)

// 4. create the object
cudaTextureObject_t tex = 0;
CUDA_CHECK(cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL));

מרכז כל טקסל יושב בקואורדינטה שלמה-ועוד-חצי, ולכן כדי לדגום את הפיקסל (px, py) פונים ל-(px + 0.5f, py + 0.5f). הנה kernel של הגדלת תמונה (upscaling) שמנצל את האינטרפולציה בחומרה - שימו לב שאין בו שום חשבון אינטרפולציה ידני ושום בדיקת גבולות, החומרה עושה הכל:

__global__ void resample(cudaTextureObject_t tex, float* out,
                         int outW, int outH, float sx, float sy) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < outW && y < outH) {
        float u = (x + 0.5f) * sx;          // map back to the source coordinate
        float v = (y + 0.5f) * sy;
        out[y * outW + x] = tex2D<float>(tex, u, v);  // bilinear sample in one operation
    }
}

בסיום, כמו כל משאב CUDA, משחררים: cudaDestroyTextureObject(tex) ואז cudaFreeArray(cuArr). יש להוסיף שלצד ה-texture (קריאה-בלבד עם דגימה) קיים גם surface memory דרך cudaSurfaceObject_t, שמאפשר גם כתיבה למבנה ה-cudaArray הדו-ממדי - שימושי כשצריך גם לקרוא וגם לכתוב תמונה בפריסת texture, למשל בסימולציות על רשת.

דוגמה מודרנית - סטנסיל 2D דרך טקסטורה - a 2D stencil via texture

נחבר את הרעיונות בדוגמה שמראה למה texture נוח לסטנסיל דו-ממדי. סטנסיל ממוצע 5-נקודתי (הליבה של פותרי משוואות דיפרנציאליות ושל טשטוש תמונה) קורא לכל פיקסל את ארבעת שכניו. ב-global memory זה מחייב בדיקות גבול בכל כיוון וגישה בלתי-מאוחדת בציר האנכי; ב-texture הכל נעלם:

__global__ void jacobi5pt(cudaTextureObject_t tex, float* out, int W, int H) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < W && y < H) {
        float xc = x + 0.5f, yc = y + 0.5f;
        float c = tex2D<float>(tex, xc,       yc);
        float l = tex2D<float>(tex, xc - 1.f, yc);   // the clamp mode handles the left edge
        float r = tex2D<float>(tex, xc + 1.f, yc);
        float u = tex2D<float>(tex, xc,       yc - 1.f);
        float d = tex2D<float>(tex, xc,       yc + 1.f);
        out[y * W + x] = 0.2f * (c + l + r + u + d);
    }
}

ה-cudaAddressModeClamp מבטיח שכשפונים ל-xc - 1 בפיקסל השמאלי ביותר, החומרה מחזירה את פיקסל הקצה במקום לקרוס או לדרוש if. הלוקליות הדו-ממדית של cache הטקסטורה נהנית מכך ששכנים אנכיים (yc - 1, yc + 1) של threads סמוכים חופפים בcache. זו הסיבה ש-texture הייתה במשך שנים ברירת המחדל לסטנסילים ולעיבוד תמונה - אף על פי שכיום, על חומרה מודרנית, גרסת shared-memory עם tiling (כפי שראינו ב-3.3) לרוב מנצחת אותה בביצועים גולמיים. עדיין, ה-texture נותרת קלה לכתיבה, ומנצחת כשהגישה ממש בלתי-סדירה או כשצריך אינטרפולציה.

בחירת המרחב הנכון - choosing the right space

עכשיו יש לנו את כל התמונה. השאלה המעשית בכל kernel היא: איזה מרחב זיכרון לתת לכל מערך. הכלל הפשוט:

• אותו ערך לכל הthreads (מקדמים, פרמטרים, טבלה קטנה שנקראת אחיד) -> constant memory.
• קריאה-בלבד עם שימוש-חוזר מרחבי, גישה דו-ממדית או אינטרפולציה -> texture (או cache קריאה-בלבד לגישה קווית).
• קריאה-בלבד קווית, סתם רמז cache -> const __restrict__ / __ldg (cache הקריאה-בלבד).
• קריאה-כתיבה כללית -> global memory מאוחד (coalesced), עם staging דרך shared memory לשימוש-חוזר בתוך block.

הטבלה המלאה, כולל מתי כל מרחב מזיק:

מרחב	תבנית אידיאלית	קיבולת	מנגנון	מתי מזיק
constant	אותה כתובת לכל ה-warp (broadcast)	64 KiB	הcache קבוע per-SM	כתובות מתפצלות -> סדרתי, עד פי 32
read-only (__ldg)	קריאה-בלבד קווית, שימוש-חוזר	כגודל ה-L1	הcache קריאה-בלבד לא-קוהרנטי	נתון שנכתב -> ערכים מיושנים
texture	לוקליות 2D, אינטרפולציה, גבולות	כגודל התמונה ב-array	cache texture + יחידת דגימה	דיוק שבור באינטרפולציה; overhead הקמה
shared	שימוש-חוזר בתוך block	256 KiB/SM ב-H100	SRAM על ה-SM, מנוהל-מתכנת	bank conflicts; scope מוגבל ל-block
global	קריאה-כתיבה כללית מאוחדת	80 GiB ב-H100	L2 + HBM3 (כ-3.35 TB/s)	גישה לא-מאוחדת -> בזבוז רוחב פס

הנקודה החינוכית: אלה לא ארבעה סוגי RAM נפרדים פיזית. constant, read-only ו-texture כולם קוראים בסופו של דבר מאותו GPU RAM כמו global - ההבדל הוא בנתיב הcache שדרכו הקריאה עוברת ובחומרה הייעודית (broadcast, יחידת דגימה) שמצטרפת אליו. הבחירה הנכונה היא התאמה בין תבנית הגישה שלכם לבין הנתיב שממוטב לה. בפרויקט 3.5 תיישמו את ההחלטות האלה על kernel אמיתי ותמדדו את ההבדל.

סיכום

• שלושת מרחבי הזיכרון המיוחדים - constant, read-only ו-texture - כולם משרתים קריאה בלבד מ-device, קוראים בסופו של דבר מאותו GPU RAM כמו global memory, ונבדלים בנתיב הcache ובחומרה הייעודית שמצטרפת אליו.
• זיכרון קבוע (__constant__, עד 64 KiB, נטען עם cudaMemcpyToSymbol) ממוטב ל-broadcast: כשכל ה-warp קורא אותה כתובת, כולם משורתים בפעולה אחת - מושלם למקדמים ולפרמטרים.
• זיכרון קבוע מזיק כשהthreads ב-warp קוראים כתובות שונות: הגישות מסתדרות בטור, עד פי 32 איטי יותר, ואז global מאוחד עדיף.
• הcache הקריאה-בלבד מנותב עם __ldg() או, עדיף, עם const __restrict__, ומופיע ב-SASS כ-LDG.E.CI; הוא נותן נתיב cache לא-קוהרנטי לנתונים שרק נקראים, ומידע אי-aliasing יקר לcompiler.
• מ-Volta והלאה (וכך ב-H100) ה-L1 וcache הקריאה-בלבד מאוחדים, ולכן הרווח מ-__ldg פחות דרמטי מבעבר, אבל const __restrict__ עדיין הרגל טוב.
• זיכרון טקסטורה דרך cudaTextureObject_t נותן בחומרה לוקליות דו-ממדית, קואורדינטות מנורמלות, מצבי גבול (clamp/wrap/mirror/border) ואינטרפולציה בי-לינארית בדיוק שבור של 8 ביטים - מצוין לעיבוד תמונה ולסטנסילים.
• ל-texture מקבילה בשם surface (cudaSurfaceObject_t) שמאפשרת גם כתיבה למבנה cudaArray דו-ממדי.
• כלל הבחירה: אותו ערך לכל הthreads -> constant; קריאה-בלבד עם מרחביות/אינטרפולציה -> texture; קריאה-בלבד קווית -> const __restrict__; קריאה-כתיבה כללית -> global מאוחד עם staging ב-shared.