לדלג לתוכן

0.3 התוכנית הראשונה חיבור וקטורים תרגול

תרגול - התוכנית הראשונה - חיבור וקטורים

בתרגול הזה תכתבו, תקמפלו ותריצו את תוכנית ה-CUDA הראשונה שלכם, ואז תרחיבו אותה צעד-צעד: תוסיפו טיפול איתן בשגיאות ותגרמו לשגיאה בכוונה כדי לראות אותה, תממשו את SAXPY, תמירו את הקוד ללולאת grid-stride, ולבסוף תמדדו זמנים ותגלו מתי ה-GPU באמת משתלם. עבדו לפי הסדר - כל תרגיל נשען על הקודם. הריצו הכל על GPU אמיתי; אם אין לכם, השתמשו ב-Google Colab (T4 חינמי) או במכונת ענן, כפי שראינו בשיעור 0.2. ה-H100 הוא כרטיס הייחוס שלנו, אבל כל תרגיל כאן ירוץ על כל GPU של NVIDIA.

הכנה

צרו קובץ vecadd.cu והעתיקו אליו את התוכנית המלאה מההרצאה (הכוללת את מקרו CUDA_CHECK, את ה-kernel בשם vecAdd, ואת main). הקומפילציה לאורך כל התרגול (החליפו את sm_90a בארכיטקטורה של הכרטיס שלכם - למשל sm_75 ל-T4):

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

כדי לברר את ה-compute capability של הכרטיס:

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

תרגיל 1 - הקלדה, קומפילציה והרצה

  1. הקלידו את vecadd.cu (אל תעתיקו בהדבקה עיוורת - הקלדה מלמדת).
  2. קמפלו עם שורת ה-nvcc שלמעלה.
  3. הריצו את ./vecadd וּודאו שהפלט הוא max error = 0.000000 ושתי שורות הפלט של c[0] ו-c[1048575].
  4. שנו את n ל-1 << 25 (כ-33.5 מיליון איברים), קמפלו והריצו שוב. ודאו שהתוצאה עדיין נכונה.

רמז: 1 << 25 דורש שלושה מערכים של 128 MiB כל אחד על ה-GPU. אם cudaMalloc נכשל בגלל חוסר זיכרון, מקרו ה-CUDA_CHECK ידפיס out of memory עם הקובץ והשורה - זה בדיוק מה שהוא נועד לתפוס.

תרגיל 2 - שגיאה מכוונת וטיפול בה

  1. ודאו שהתוכנית שלכם עוטפת את הlaunch בשתי הבדיקות: CUDA_CHECK(cudaGetLastError()) ואז CUDA_CHECK(cudaDeviceSynchronize()).
  2. גרמו לשגיאת launch בכוונה: שנו את blockSize מ-256 ל-2048 (מעל המקסימום של 1024 threads ל-block). קמפלו והריצו, ורשמו את ההודעה המדויקת שקיבלתם.
  3. החזירו את blockSize ל-256. עכשיו גרמו לשגיאת ריצה במקום launch: הסירו את שומר הגבולות if (i < n) (השאירו רק את c[i] = a[i] + b[i];) ושנו את numBlocks כך שיושק יותר מ-n threads - למשל הוסיפו numBlocks += 100;. קמפלו, הריצו, וראו איזו שגיאה נתפסת ובאיזו בדיקה.
  4. הסבירו לעצמכם: איזו שגיאה נתפסה על ידי cudaGetLastError ואיזו על ידי cudaDeviceSynchronize, ולמה.

רמז: תצורת הרצה לא חוקית מתגלה מיד בlaunch (cudaGetLastError מחזיר invalid configuration argument). גישה מחוץ לתחום מתגלה רק כשה-kernel רץ בפועל, ולכן היא נתפסת ב-cudaDeviceSynchronize (an illegal memory access was encountered). זכרו: שגיאת ריצה כזו "דביקה" - אחריה כל קריאת CUDA תיכשל.

תרגיל 3 - מימוש SAXPY

הפעולה SAXPY (Single-precision A times X Plus Y) מחשבת y = a*x + y, כאשר a סקלר ו-x, y וקטורים. זו אבן היסוד של אלגברה לינארית ומופיעה ב-cuBLAS.

  1. כתבו kernel חדש saxpy(int n, float a, const float* x, float* y). שימו לב ש-y הוא גם קלט וגם פלט - הוא נקרא ונכתב באותו איבר.
  2. בתוך main, אתחלו x[i] = 1.0f, y[i] = 2.0f, ובחרו a = 3.0f.
  3. השיקו את ה-kernel, העתיקו את התוצאה חזרה, ואמתו שכל איבר שווה ל-3*1 + 2 = 5.
  4. מהו ה-max error שאתם מצפים לראות?

רמז: מבנה ה-kernel זהה ל-vecAdd - חשבו את i בדיוק אותו הדבר, השתמשו בשומר גבולות if (i < n), והשורה היחידה בפנים היא y[i] = a * x[i] + y[i];. הסקלר a עובר by-value בתצורת הארגומנטים הרגילה, בלי cudaMalloc - רק מצביעים דורשים זיכרון התקן.

תרגיל 4 - המרה ללולאת grid-stride

  1. כתבו גרסה vecAddStride של ה-kernel שמשתמשת בלולאת grid-stride (ראו את התבנית בהרצאה).
  2. הריצו אותה עם grid קטן במכוון מהנתונים: קבעו numBlocks = 128 קבוע (במקום (n + blockSize - 1) / blockSize), עם blockSize = 256 ו-n = 1 << 20. זהו grid של 32,768 threads בלבד לעומת מיליון איברים.
  3. ודאו שהתוצאה עדיין נכונה (max error = 0) למרות שיש הרבה פחות threads מאיברים.
  4. חשבו: כמה איברים מטפל בהם כל thread בממוצע במקרה הזה?

רמז: stride = blockDim.x * gridDim.x = 256 * 128 = 32768. כל thread מטפל בערך n / stride = 1048576 / 32768 = 32 איברים. אם התוצאה שגויה, בדקו שהצעד בלולאה הוא i += stride (ולא i++) ושה-stride מחושב מ-gridDim.x ולא מקבוע.

תרגיל 5 - מדידת זמנים: CPU מול GPU

  1. כתבו לולאת CPU פשוטה שמחשבת h_ref[i] = h_a[i] + h_b[i] ומדדו את זמנה עם std::chrono (או clock()).
  2. מדדו את זמן ה-kernel ב-GPU עם אירועי CUDA - cudaEvent:
cudaEvent_t start, stop;
CUDA_CHECK(cudaEventCreate(&start));
CUDA_CHECK(cudaEventCreate(&stop));
CUDA_CHECK(cudaEventRecord(start));
vecAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);
CUDA_CHECK(cudaEventRecord(stop));
CUDA_CHECK(cudaEventSynchronize(stop));
float ms = 0.0f;
CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
  1. מדדו בנפרד גם את זמן ההעברות (cudaMemcpy H2D ו-D2H) עם אותם אירועים.
  2. הריצו עבור n הולך וגדל: 1 << 10, 1 << 16, 1 << 20, 1 << 24, 1 << 27. בנו טבלה של זמן CPU, זמן kernel בלבד, וזמן kernel + העברות.
  3. ענו: באיזה n מנצח ה-kernel לבדו את ה-CPU? ומה קורה למסקנה כשמוסיפים את זמן ההעברות?

רמז: מדדו את ה-kernel לבד (בלי ההעברות) כי אירוע CUDA מתזמן את הזרם (stream) של ה-GPU. הריצו את ה-kernel פעם אחת ל"חימום" לפני המדידה (הקומפילציה של PTX ל-SASS והקצאות ראשונות מוסיפות overhead חד-פעמית). זכרו ש-cudaEventElapsedTime מחזיר מילישניות כ-float. תראו שזמן ה-kernel זעום גם ל-n גדול, אבל זמן ההעברות על ה-PCIe שולט - זו הסיבה שחיבור וקטורים לבדו כמעט אף פעם לא "שווה" העברה ל-GPU.

תרגיל 6 (בונוס) - בדיקת רוחב פס

  1. עבור ההרצה של n = 1 << 27 (כ-134 מיליון איברים), חשבו את התעבורה מול הזיכרון: כל איבר קורא שני float וכותב אחד, כלומר 3 * 4 = 12 בתים לאיבר.
  2. מתוך זמן ה-kernel שמדדתם, חשבו את רוחב הפס האפקטיבי ב-GB/s: (12 * n) / (זמן בשניות) / 1e9.
  3. השוו לרוחב הפס התיאורטי של הכרטיס שלכם (ב-H100 SXM כ-3.35 TB/s; אפשר לברר עם nvidia-smi -q או deviceQuery). איזה אחוז ניצלתם?

רמז: רוחב פס אפקטיבי גבוה (עשרות אחוזים מהתיאורטי) הוא סימן שהגישות מאוחדות (coalesced) והחישוב memory-bound כפי שצפינו. אם קיבלתם אחוז נמוך מאוד, ודאו שמדדתם את ה-kernel לבד ולא כללתם את זמן ההעברות. את המסגרת המלאה לניתוח - מודל ה-Roofline וarithmetic intensity - נלמד בפרק על ביצועים.