לדלג לתוכן

0.2 סביבת עבודה תרגול

תרגול - סביבת עבודה

בתרגול הזה תוודאו שסביבת ה-CUDA שלכם עובדת מקצה לקצה - מהדרייבר ועד ל-kernel שרץ באמת על ה-GPU. תקראו את מצב הכרטיס, תאמתו שהדרייבר וה-Toolkit מותקנים, תכתבו תוכנית שחוקרת את הכרטיס, ותקמפלו ותריצו את ה-kernel הראשון שלכם. עבדו לפי הסדר: כל תרגיל בונה על הקודם. אפשר לעשות את הכל על כרטיס NVIDIA מקומי, או בחינם ב-Google Colab. אם אתם ב-Colab, החליפו -arch=sm_90 ב--arch=sm_75 (זהו ה-compute capability של ה-T4), והוסיפו סימן קריאה (!) לפני כל פקודת shell בתא.

הכנה - שני קובצי מקור

לאורך התרגול נשתמש בשני קבצים. צרו אותם עכשיו. הראשון, devinfo.cu, חוקר את הכרטיס:

// devinfo.cu
#include <cstdio>
#include <cstdlib>
#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)

int main() {
    // fill in during exercise 3
    return 0;
}

השני, hello.cu, נשלים אותו בתרגיל 4. קמפלו תמיד עם ציון הארכיטקטורה, למשל nvcc -arch=sm_90 devinfo.cu -o devinfo.

תרגיל 1 - קריאת מצב הכרטיס

לפני שכותבים שורת קוד אחת, מסתכלים על הכרטיס.

  1. הריצו nvidia-smi וודאו שהיא רצה בלי שגיאה (אם לא - הדרייבר לא מותקן).
  2. מהשורה העליונה, רשמו את גרסת הדרייבר (Driver Version) ואת גרסת ה-CUDA המקסימלית (CUDA Version). הסבירו במילים שלכם למה שני המספרים האלה שונים ומה כל אחד אומר.
  3. מהטבלה, זהו וכתבו עבור הכרטיס שלכם: הדגם (Name), הטמפרטורה (Temp), מצב הביצועים (Perf), צריכת ההספק מול התקרה (Pwr:Usage/Cap), וכמה זיכרון פנוי (Memory-Usage).
  4. הריצו nvidia-smi -q ומצאו בפלט את סף ההאטה התרמי (GPU Slowdown Temp) ואת תדר ה-SM הנוכחי. השוו את תדר ה-SM בבטלה לתדר ה-boost המלא של הכרטיס.

רמז: השדה CUDA Version בכותרת הוא התקרה שהדרייבר תומך בה, לא בהכרח ה-Toolkit המותקן. השתמשו ב-nvidia-smi -q | less וחפשו את הקטעים Temperature ו-Clocks.

תרגיל 2 - הוכחה שהדרייבר וה-Toolkit מותקנים

עכשיו נאמת את שני חצאי המחסנית בנפרד.

  1. הריצו nvcc --version ורשמו את גרסת ה-Toolkit (השורה release).
  2. מצאו את libcuda.so (מגיעה מהדרייבר) ואת libcudart.so (מגיעה מה-Toolkit) במערכת. השתמשו ב-find / -name 'libcuda.so*' 2>/dev/null ו-find / -name 'libcudart.so*' 2>/dev/null.
  3. הסבירו: אם libcuda.so הייתה חסרה, מה זה היה אומר? ואם רק nvcc היה חסר?
  4. אחרי שתקמפלו את devinfo בתרגיל 3, הריצו ldd ./devinfo וזהו מולן מ-הספריות ה-CUDA הבינארי נקשר בזמן ריצה.

רמז: libcuda.so מגיעה מחבילת הדרייבר ו-libcudart.so מה-Toolkit - זו הסיבה שהיעדר libcuda.so מצביע על דרייבר שבור, ולא על Toolkit חסר. שימו לב שלפעמים libcudart נקשרת סטטית ואז לא תופיע ב-ldd.

תרגיל 3 - חקירת הכרטיס מתוך קוד

השלימו את devinfo.cu כך שתדפיס את מאפייני הכרטיס.

  1. בתוך main, קראו ל-cudaGetDeviceCount (עטוף ב-CUDA_CHECK) ולולאה על כל device.
  2. לכל device קראו ל-cudaGetDeviceProperties, והדפיסו לפחות: השם, ה-compute capability (major.minor), מספר ה-SM (multiProcessorCount), ה-warpSize, maxThreadsPerBlock, sharedMemPerBlock, וגודל הזיכרון הגלובלי (totalGlobalMem).
  3. קמפלו והריצו. רשמו עבור הכרטיס שלכם: כמה SM, איזה compute capability, כמה HBM (בג'יגה), ומה ה-warp size.
  4. השתמשו במספרים כדי לחשב את שיא ה-FP32 של הכרטיס: SM_count x 128 x 2 x boost_clock_GHz (מניחים 128 ליבות FP32 ל-SM בדורות המודרניים). השוו לתוצאה של H100 (66.9 TFLOPS).

רמז: totalGlobalMem הוא ב-bytes מסוג size_t - חלקו ב-1024.0*1024.0*1024.0 בשביל GiB, והשתמשו ב-%zu להדפסת שדות size_t. את תדר ה-boost תשלפו מ-clockRate (ב-kHz) או מ-nvidia-smi -q.

תרגיל 4 - ה-kernel הראשון

הגיע הרגע לכתוב קוד שרץ על ה-GPU.

  1. מלאו את hello.cu עם kernel המסומן __global__ שמדפיס, מכל thread, את האינדקס הגלובלי שלו (blockIdx.x * blockDim.x + threadIdx.x), את מספר ה-block ואת ה-lane (threadIdx.x).
  2. ב-main, השיקו את ה-kernel עם <<<2, 4>>> (2 blocks, 4 threads בכל אחד) והוסיפו cudaDeviceSynchronize() אחריו.
  3. קמפלו והריצו. ודאו שמודפסות בדיוק 8 שורות.
  4. הריצו את התוכנית כמה פעמים ברצף. שימו לב האם סדר השורות זהה בכל הרצה. הסבירו מה שאתם רואים.
  5. הסירו בכוונה את cudaDeviceSynchronize(), קמפלו והריצו שוב. מה קורה לפלט, ולמה?

רמז: kernel תמיד מחזיר void - הוא לא "מחזיר" ערך אלא כותב (כאן, מדפיס). בלי cudaDeviceSynchronize() הhost יכול לצאת לפני שה-GPU רוקן את חוצץ ה-printf, אז הפלט עלול להיעלם. הסדר בין blocks לעולם אינו מובטח.

תרגיל 5 - תפיסת שגיאה עם CUDA_CHECK

המקרו CUDA_CHECK שווה רק אם הוא באמת תופס שגיאות. נגרום לו לירות.

  1. בתוכנית נפרדת (או ב-devinfo.cu זמנית), קראו ל-cudaGetDeviceProperties עם אינדקס לא חוקי, למשל 999, כשהוא עטוף ב-CUDA_CHECK.
  2. קמפלו והריצו. ודאו שהתוכנית יוצאת עם הודעה שמכילה את שם הקובץ, מספר השורה, וטקסט השגיאה.
  3. רשמו את מחרוזת השגיאה המדויקת שקיבלתם.

רמז: cudaGetErrorString מתרגם cudaError_t לטקסט קריא. אינדקס device לא חוקי מחזיר cudaErrorInvalidDevice.

תרגיל 6 (בונוס) - הכל בענן

אם עשיתם הכל מקומית, עשו את זה עכשיו ב-Google Colab (או ב-Modal); ואם עבדתם בענם - עשו את זה מקומית.

  1. פתחו notebook ב-Colab, בחרו runtime מסוג GPU (Runtime -> Change runtime type -> GPU).
  2. הריצו !nvidia-smi וזהו איזה כרטיס קיבלתם ומה ה-compute capability שלו (בדרך כלל T4, שהוא 7.5).
  3. העלו או צרו את devinfo.cu ו-hello.cu, קמפלו עם !nvcc -arch=sm_75 ... והריצו.
  4. השוו את מספר ה-SM ושיא ה-FP32 של הכרטיס בענן לאלה של H100. פי כמה H100 חזק יותר?

רמז: ב-Colab כל פקודת shell מתחילה ב-!. כדי לכתוב קובץ מתא, אפשר להשתמש ב-%%writefile devinfo.cu בתחילת התא. ל-T4 יש 40 SM; חשבו לבד את שיא ה-FP32 שלו והשוו.