הבלוג של ארתיום
בלוג על לינוקס, תוכנה חופשית, מוזיקה, סלסה, ומה לא!
מאמרים בנושא פיתוח.
התקדמות חשובה בתמיכה ב־OpenCL ב־pytorch.
רקע
היום pytorch היא אחת התשתיות המובילות בעולם למידה עמוקה. יש לה יתרונות רבות, אבל מבחינת המפתח זה קוד איכותי ותיעוד טוב. הקוד כתוב בצורה מאוד מודרנית עם שימוש נכון ביכולות C++ מה שמאוד מקל על הפיתוח. אני עובד בתקופה האחרונה על פיתוח מנוע עבור pytorch מבוסס OpenCL כחלופה ל־cuda.
כבר כתבתי בעבר על חשיבות התמיכה ב־OpenCL.
אבל בכל זאת אזכיר כמה נקודות מבחינת קהילת תוכנה חופשית וקוד פתוח:
- אנחנו זקוקים בתמיכה חוצת פלטפורמה בכרטיסי מסך מיצרנים שונים כמו AMD, Intel וכמובן nVidia.
- אנחנו זקוקים למימוש האלגוריתמים המרכזיים כקוד פתוח הזמין לכל (ולא כקופסה סגורה ש־nVidia נותנת)
- אנחנו רוצים לעבוד עם סטנדרטים פתוחים וזמינים כמו OpenCL ולא מימושים ספציפיים של יצרן (כמו cuda).
אז מה חדש? קלות אינטגרציה!
עם שחרור גרסה 1.13 של pytorch חל שיפור משמעותי ב־out-of-tree-backend. עכשיו הוספת מנוע אימון מבוסס OpenCL היא פשוטה מאוד ולמעשה שאלה של מספר דקות, אפילו בוונידוס העניין יחסית פשוט. המשמעות שאפשר להשתמש במנוע OpenCL בקלות רבה הן בלינוקס והן בווינדוס.
מה עם הביצועים? אם משווים מול גרסת cuda/cudnn על אותו ה־gpu מדובר בין 50 ל־70 אחוז ביצועי cuda באימון ובין כ־60 ל־80 באבלואציה (תלוי ברשת כמובן).
למרות שהמנוע עדיין ניסיוני וחסרים בו לא מעט פעולות הוא נבדק בהצלחה על עשרות רשתות כמו resnet, mobilenet ורבות אחרות.
המנוע עצמו מבוסס על ספריית dlprimitives שאני מפתח במקביל והיא חלופה ל־cuDNN על בסיס OpenCL וגם מנוע חיזוי שעובד עם מודלים בפורמט ONNX - שזה נושא גדול בפני עצמו.
מה המשמעות של זה?
משתמשי AMD יכולים לאמן רשתות. הם לא מוגבלים למספר מצומצם של דגמים ש־rocm תומך בהם או לשימוש בלינוקס בלבד. התמיכה היא גורפת מ־APUים ישנים כמו Stoney Ridge ועד ל־RDNA 2 וגם זה עובד על "חלונות" למי שמעוניין.
זו הייתה משימה כמעט ובלי אפשרית עד היום. עכשיו זה במרחק מספר פקודות
תשתית אימון היא קוד פתוח לגמרי גם אם עובדים עם nVidia (טוב חוץ מהדרייבר שלהם)
- כל מה שצריך זה דרייברי של OpenCL. לא צריך את כל המפלצת של cuda (מי שיצא לו להתקין לשדרג לגלות בעיות תאימות יבין אותי מידי)
מחפש עזרה...
מישהו יודע איך אפשר לבנות ולפרסם whl לפלטפורמות שונות? רצוי איזה שירות ענן שיעשה זאת? כדי שזה יהיה ממש במרחק של pip install :-)
מעשה בשני NaNים
לאחרונה ניסיתי להריץ אימון של GAN על pytorch עם תמיכה ב־dlprimitives. אחד דברים הלא נעימים באימון של GANים באופן כללי זה שהם לא ממש יציבים ומתבדרים בקלות.
שמתי לב שבגרסת cuda הוא רץ ללא דופי ובגרסה שלי הוא נתקע. נתקע על ביצוע backpropogation ב־convolution. אחד ההבדלים העיקריים באגלוריתם בהשוואה לשאר היה שימוש בפעולות אטומיות לחישוב סכום (טריק מאוד נפוץ במימוש קונבולוציה)
אחרי זמן מה הגעתי למסקנה שהחישובים מגיעים ל־NaN באיזהו מקום ואז הכל נתקע. לא הבנתי למה פעולת חיבור אטומית פשוטה נתקעת. בכל מקרה איתרתי באג האחר שהביא לחישוב ה־NaN והכל הסתדר. אבל בכל זאת נושא התקיעה הטריד אותי.
כתבתי שאלה ב־Stackoverflow העתקתי קטע קוד... ואז נפל האסימון
float oldv = *ptr;
for(;;) {
float newv = oldv + v;
float prev = as_float(atomic_cmpxchg((__global volatile int *)(ptr),as_int(oldv),as_int(newv)));
if(prev == oldv)
return;
oldv = prev;
}
קחו לכם כמה דקות.
פעולת comxchg ב־OpenCL עובדת רק על int. לכן הייתי צריך לעשות bit-casting ל־int ובחזרה (as_float ו־as_int בדיוק עושים את זה). ואז תנאי הבדיקה prev==old ביצעתי ב־float במקום בשלמים.
כך שאם הערכים שווים אז ההחלפה הצליחה. אבל מה שכחתי? NaN == NaN תמיד נותן false! ולכן תקעתי בלולאה אינסופית כי התנאי לעולם לא ייתקיים. החלפתי לבדיקה בערכים שלמים (קרי ייצוג בינארי) והכל עבד חלק.
מסקנה NaN הוא טריקי... יש לכבדו
התחלתי להתקדם לאימון בקוד פתוח: pytorch עם תמיכה ב־OpenCL
תקציר: הצלחתי לבצע inference של AlexNet ב־pytorch ב־OpenCL. הביצועים זהים לפעלה ישירה של dlprimitives.
הדרך עוד ארוכה אבל פחות קצת יותר ברורה מה לעשות. המאמר המלא באנגלית בבלוג הפיתוח:
http://blog.dlprimitives.org/post/5
השוואה בין תפוחים ירוקים ואדומים
כשמודבר ב־Deep Learning זה מאוד קשה להשוות בין GPU של החברות המובילות AMD ו־NVidia. בניגוד למשחקי מחשב שנותנים לך מדדים ברורים על עלות מול תועלת, תחום DL נשלט באופן בלעדי ע"י NVidia. גם אם קיימים פתרונות של AMD הם לא תמיד עובדים. למשל כרטיסי RDNA/RDNA2 עדיין לא נתמכים ע"י AMD לטובת Deep Learning - והם בעצם הכרטיסים הזמינים היחידים היום בשוק.
ובכן כחלק מפרויקט DLPrimitives עשיתי השוואה כזו:
http://blog.dlprimitives.org/post/1
וכן גם השקתי בלוג חדש לטובת עדכונים על הפרויקט.
רשתות נוירונים בקוד פתוח... להפשיל שרוולים
כתבתי בעבר על המצב העגום של תחום ה־deep learning בקוד פתוח - שלמעשה לא קיים. אחרי ש־Google בפועל הרגו את ה־plaidml עם keras והפיתוח של Caffe הופסק אז נוצר המצב בו אין כל דרך לאמן רשתות בעזרת פלטפורמה פתוחה - OpenCL.
יש סיבות טובות לעבוד עם OpenCL מעבר לשמירה על הקוד הפתוח. למשל לפתח תוכנה שתעבוד על כל כרטיס גרפי סביר ובכל מערכת הפעלה - בלי להסתבך.
אז הרמתי את הכפפה: https://github.com/artyom-beilis/dlprimitives
זהו פרויקט חדש בשם DLPrimitives שאמור לתת מענה לסוגיה. הוא אמור לספק ספריה בסגנון cudnn/miopen שמממשת את הפעולות הבסיסיות של Deep-Learning וגם לספק כלים ל־inference. בנוסף, הרעיון הוא להתחבר כ־backend לאחד ה־deep learning frameworks העדכניים כמו pytorch, tensorflow או mxnet.
התהליך הוא איטי וקשה. אומנם המתמטיקה היא לא מסובכת וכתיבה ל־GPU היא בסה"כ לא עניין מסובך. אבל אם רוצים להגיע לביצועים טובים הסיפור הוא מעט שונה. עם זה, התוצאות כבר כאן.
לחסרי סבלנות - הצלחתי להגיע ל־150%-200% של ביצועי caffe-opencl ו־plaidml על פלטפורמת amd ו־nvidia ולהגיע לכ־50% עד 70% של ביצועי המימושים הספציפיים שלהם על בסיס cudnn/miopen.
כל התוצאות:
https://github.com/artyom-beilis/dlprimitives/blob/master/docs/summary.md
סיכום לעצלנים - ממוצע על 5 רשתות נפוצות alexnet, resnet18, resnet50, vgg, mobilenet:
| GPU | Batch | Train, Cuda/HIP | Test, Cuda/HIP | Train, Plaidml/Caffe | Test, Plaidml/Caffe |
|---|---|---|---|---|---|
| gtx960 | 16 | 51% | 60.73% | 171% | 167.33% |
| gtx960 | 8 | 59% | 72.03% | 187% | 155.25% |
| gtx1080 | 16 | 42% | 41.34% | 207% | 137.52% |
| rtx2060s | 16 | 49% | 57.53% | 211% | 149.48% |
| rx560 | 16 | 53% | 56.82% | 153% | 115.63% |
| rx560 | 8 | 55% | 54.19% | 172% | 122.64% |
| intel-hd530 | 8 | 109% | 66.12% |
אומנם זו התחלה אבל כבר התחלה טובה!
