הבלוג של ארתיום :: בינה מלאכותית http://artyom.cppcms.com/ בלוג על לינוקס, תוכנה חופשית, מוזיקה, סלסה, ומה לא! התקדמות חשובה בתמיכה ב־OpenCL ב־pytorch. http://artyom.cppcms.com/post/337 http://artyom.cppcms.com/post/337 <div style="direction:rtl"> <h2>רקע</h2> <p>היום pytorch היא אחת התשתיות המובילות בעולם למידה עמוקה. יש לה יתרונות רבות, אבל מבחינת המפתח זה קוד איכותי ותיעוד טוב. הקוד כתוב בצורה מאוד מודרנית עם שימוש נכון ביכולות C++‎ מה שמאוד מקל על הפיתוח. אני עובד בתקופה האחרונה על פיתוח מנוע עבור pytorch מבוסס OpenCL כחלופה ל־cuda.</p> <p>כבר כתבתי <a href="http://blog.dlprimitives.org/post/2">בעבר</a> על חשיבות התמיכה ב־OpenCL.</p> <p>אבל בכל זאת אזכיר כמה נקודות מבחינת קהילת תוכנה חופשית וקוד פתוח:</p> <ol> <li>אנחנו זקוקים בתמיכה חוצת פלטפורמה בכרטיסי מסך מיצרנים שונים כמו AMD, Intel וכמובן nVidia.</li> <li>אנחנו זקוקים למימוש האלגוריתמים המרכזיים כקוד פתוח הזמין לכל (ולא כקופסה סגורה ש־nVidia נותנת)</li> <li>אנחנו רוצים לעבוד עם סטנדרטים פתוחים וזמינים כמו OpenCL ולא מימושים ספציפיים של יצרן (כמו cuda).</li> </ol> <p><a href="https://github.com/artyom-beilis/pytorch_dlprim">הפרוייקט ב־github‏</a></p> <h2>אז מה חדש? קלות אינטגרציה!</h2> <p>עם שחרור גרסה 1.13 של pytorch חל שיפור משמעותי ב־out-of-tree-backend. עכשיו הוספת מנוע אימון מבוסס OpenCL היא פשוטה מאוד ולמעשה שאלה של מספר דקות, אפילו בוונידוס העניין יחסית פשוט. המשמעות שאפשר להשתמש במנוע OpenCL בקלות רבה הן בלינוקס והן בווינדוס.</p> <p>מה עם הביצועים? אם משווים מול גרסת cuda/cudnn על אותו ה־gpu מדובר בין 50 ל־70 אחוז ביצועי cuda באימון ובין כ־60 ל־80 באבלואציה (תלוי ברשת כמובן).</p> <p>למרות שהמנוע עדיין ניסיוני וחסרים בו לא מעט פעולות הוא נבדק בהצלחה על עשרות רשתות כמו resnet, mobilenet ורבות אחרות.</p> <p>המנוע עצמו מבוסס על ספריית <a href="https://github.com/artyom-beilis/dlprimitives">dlprimitives‏</a> שאני מפתח במקביל והיא חלופה ל־cuDNN על בסיס OpenCL וגם מנוע חיזוי שעובד עם מודלים בפורמט ONNX - שזה נושא גדול בפני עצמו.</p> <h2>מה המשמעות של זה?</h2> <ul> <li><p>משתמשי AMD יכולים לאמן רשתות. הם לא מוגבלים למספר מצומצם של דגמים ש־rocm תומך בהם או לשימוש בלינוקס בלבד. התמיכה היא גורפת מ־APUים ישנים כמו Stoney Ridge ועד ל־RDNA 2 וגם זה עובד על "חלונות" למי שמעוניין.</p> <p> זו הייתה משימה כמעט ובלי אפשרית עד היום. עכשיו זה במרחק מספר פקודות</p></li> <li><p>תשתית אימון היא קוד פתוח לגמרי גם אם עובדים עם nVidia (טוב חוץ מהדרייבר שלהם)</p></li> <li>כל מה שצריך זה דרייברי של OpenCL. לא צריך את כל המפלצת של cuda (מי שיצא לו להתקין לשדרג לגלות בעיות תאימות יבין אותי מידי)</li> </ul> <h3>מחפש עזרה...</h3> <p>מישהו יודע איך אפשר לבנות ולפרסם whl לפלטפורמות שונות? רצוי איזה שירות ענן שיעשה זאת? כדי שזה יהיה ממש במרחק של pip install <code>:-)</code></p> </div> מעשה בשני NaNים http://artyom.cppcms.com/post/336 http://artyom.cppcms.com/post/336 <div style="direction:rtl"> <p>לאחרונה ניסיתי להריץ אימון של GAN על pytorch עם תמיכה ב־dlprimitives. אחד דברים הלא נעימים באימון של GANים באופן כללי זה שהם לא ממש יציבים ומתבדרים בקלות.</p> <p>שמתי לב שבגרסת cuda הוא רץ ללא דופי ובגרסה שלי הוא נתקע. נתקע על ביצוע backpropogation ב־convolution. אחד ההבדלים העיקריים באגלוריתם בהשוואה לשאר היה שימוש בפעולות אטומיות לחישוב סכום (טריק מאוד נפוץ במימוש קונבולוציה)</p> <p>אחרי זמן מה הגעתי למסקנה שהחישובים מגיעים ל־NaN באיזהו מקום ואז הכל נתקע. לא הבנתי למה פעולת חיבור אטומית פשוטה נתקעת. בכל מקרה איתרתי באג האחר שהביא לחישוב ה־NaN והכל הסתדר. אבל בכל זאת נושא התקיעה הטריד אותי.</p> <p>כתבתי שאלה ב־Stackoverflow העתקתי קטע קוד... ואז נפל האסימון</p> <pre><code>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; } </code></pre> <p>קחו לכם כמה דקות.</p> <p>פעולת comxchg ב־OpenCL עובדת רק על int. לכן הייתי צריך לעשות bit-casting ל־int ובחזרה (<code>as_float</code> ו־<code>as_int</code> בדיוק עושים את זה). ואז תנאי הבדיקה <code>prev==old</code> ביצעתי ב־float במקום בשלמים.</p> <p>כך שאם הערכים שווים אז ההחלפה הצליחה. אבל מה שכחתי? NaN == NaN תמיד נותן false! ולכן תקעתי בלולאה אינסופית כי התנאי לעולם לא ייתקיים. החלפתי לבדיקה בערכים שלמים (קרי ייצוג בינארי) והכל עבד חלק.</p> <p>מסקנה NaN הוא טריקי... יש לכבדו</p> </div> עדכוני dlprimitives http://artyom.cppcms.com/post/334 http://artyom.cppcms.com/post/334 <div style="direction:rtl"> <p>מספר עדכונים:</p> <ul> <li><p>התקדמות יפה עם pytorch - בוצעה ולידציה של מרבית הרשתות של סיווג הנמצאות ה־torchvision:</p> <ul> <li><code>alexnet</code></li> <li><code>resnet18</code></li> <li><code>resnet50</code></li> <li><code>vgg16</code></li> <li><code>densenet161</code></li> <li><code>googlenet</code></li> <li><code>squeezenet1_0</code></li> <li><code>inception_v3</code></li> <li><code>shufflenet_v2_x1_0</code></li> <li><code>mobilenet_v2</code></li> <li><code>mobilenet_v3_large</code></li> <li><code>mobilenet_v3_small</code></li> <li><code>resnext50_32x4d</code></li> <li><code>wide_resnet50_2</code></li> <li><code>mnasnet1_0</code></li> <li><code>efficientnet_b0</code></li> <li><code>efficientnet_b4</code></li> <li><code>regnet_y_400mf</code></li> </ul> <p> מה שאומר שניתן לאמן עכשיו הרבה סוגי רשתות. על הדרך מצאתי לא מעט באגים ותיקנתי אותם</p></li> <li><p>ניסיתי לעשות אינטגרציה עם OneDNN של אינטל (מאין cudnn ל־GPU שלהם) רק כדי לגלות ש<a href="https://github.com/oneapi-src/oneDNN/issues/1194">הביצועים שלהם בפורמט NCHW גרועים</a>. כיוון שרוב התשתיות עובדות עם הפורמט הזה pytorch, caffe, mxnet ועוד אז OneDNN לא ממש רלוונטי בינתיים. אחזור לשם כשיקרה אחד מהשניים:</p> <ul> <li>אינטל יתקנו את הביצועים עבור הפורמט הנפוץ</li> <li>אני אפתח תמיכה ב־NHWC לטובת TensorFlow בו זה פורמט ברירת מחדל</li> </ul> </li> </ul> <p>המשך יבוא</p> </div> התחלתי להתקדם לאימון בקוד פתוח: pytorch עם תמיכה ב־OpenCL http://artyom.cppcms.com/post/333 http://artyom.cppcms.com/post/333 <div style="direction:rtl"> <p>תקציר: הצלחתי לבצע inference של AlexNet ב־pytorch ב־OpenCL. הביצועים זהים לפעלה ישירה של dlprimitives.</p> <p>הדרך עוד ארוכה אבל פחות קצת יותר ברורה מה לעשות. המאמר המלא באנגלית בבלוג הפיתוח:</p> <p><a href="http://blog.dlprimitives.org/post/5">http://blog.dlprimitives.org/post/5</a></p> </div> מדוע אנחנו זקוקים ל־deep-learning מבוסס OpenCL? http://artyom.cppcms.com/post/332 http://artyom.cppcms.com/post/332 <div style="direction:rtl"> <p>במאמר חדש <a href="http://blog.dlprimitives.org/post/2">http://blog.dlprimitives.org/post/2</a> אני סוקר את הסיבות והצורך בהקמת תשתית למידה חישובית מבוססת OpenCL. אני מתייחס כאן לסיבות שהן לאו דווקא סיבות "אידואולוגיות" כמו שימוש בקוד פתוח אלא גם מתייחס לנושאים טכניים ומהותיים:</p> <ul> <li>מחקר ואלגוריתמים</li> <li>וניהול פרויקטים לטווח רחוק</li> <li>שיפור מוצר ע"י תחרותיות</li> </ul> </div> השוואה בין תפוחים ירוקים ואדומים http://artyom.cppcms.com/post/331 http://artyom.cppcms.com/post/331 <div style="direction:rtl"> <p>כשמודבר ב־Deep Learning זה מאוד קשה להשוות בין GPU של החברות המובילות AMD ו־NVidia. בניגוד למשחקי מחשב שנותנים לך מדדים ברורים על עלות מול תועלת, תחום DL נשלט באופן בלעדי ע"י NVidia. גם אם קיימים פתרונות של AMD הם לא תמיד עובדים. למשל כרטיסי RDNA/RDNA2 עדיין לא נתמכים ע"י AMD לטובת Deep Learning - והם בעצם הכרטיסים הזמינים היחידים היום בשוק.</p> <p>ובכן כחלק מפרויקט <a href="https://github.com/artyom-beilis/dlprimitives">DLPrimitives</a> עשיתי השוואה כזו:</p> <p><a href="http://blog.dlprimitives.org/post/1">http://blog.dlprimitives.org/post/1</a></p> <p>וכן גם השקתי בלוג חדש לטובת עדכונים על הפרויקט.</p> </div> רשתות נוירונים בקוד פתוח... להפשיל שרוולים http://artyom.cppcms.com/post/330 http://artyom.cppcms.com/post/330 <div style="direction:rtl"> <p>כתבתי בעבר על המצב העגום של <a href="http://artyom.cppcms.com/post/328">תחום ה־deep learning בקוד פתוח</a> - שלמעשה לא קיים. אחרי ש־Google בפועל <a href="https://github.com/plaidml/plaidml/issues/586">הרגו</a> את ה־plaidml עם keras והפיתוח של Caffe הופסק אז נוצר המצב בו אין כל דרך לאמן רשתות בעזרת פלטפורמה פתוחה - OpenCL.</p> <p>יש סיבות טובות לעבוד עם OpenCL מעבר לשמירה על הקוד הפתוח. למשל לפתח תוכנה שתעבוד על כל כרטיס גרפי סביר ובכל מערכת הפעלה - בלי להסתבך.</p> <p>אז הרמתי את הכפפה: <a href="https://github.com/artyom-beilis/dlprimitives">https://github.com/artyom-beilis/dlprimitives</a></p> <p>זהו פרויקט חדש בשם DLPrimitives שאמור לתת מענה לסוגיה. הוא אמור לספק ספריה בסגנון cudnn/miopen שמממשת את הפעולות הבסיסיות של Deep-Learning וגם לספק כלים ל־inference. בנוסף, הרעיון הוא להתחבר כ־backend לאחד ה־deep learning frameworks העדכניים כמו pytorch, tensorflow או mxnet.</p> <p>התהליך הוא איטי וקשה. אומנם המתמטיקה היא לא מסובכת וכתיבה ל־GPU היא בסה"כ לא עניין מסובך. אבל אם רוצים להגיע לביצועים טובים הסיפור הוא מעט שונה. עם זה, התוצאות כבר כאן.</p> <p>לחסרי סבלנות - הצלחתי להגיע ל־150%-200% של ביצועי caffe-opencl ו־plaidml על פלטפורמת amd ו־nvidia ולהגיע לכ־50% עד 70% של ביצועי המימושים הספציפיים שלהם על בסיס cudnn/miopen.</p> <p>כל התוצאות:</p> <p><a href="https://github.com/artyom-beilis/dlprimitives/blob/master/docs/summary.md">https://github.com/artyom-beilis/dlprimitives/blob/master/docs/summary.md</a></p> <p>סיכום לעצלנים - ממוצע על 5 רשתות נפוצות alexnet, resnet18, resnet50, vgg, mobilenet:</p> <table dir="ltr"> <thead> <tr> <th>GPU</th> <th>Batch</th> <th>Train, Cuda/HIP</th> <th>Test, Cuda/HIP</th> <th>Train, Plaidml/Caffe</th> <th>Test, Plaidml/Caffe</th> </tr> </thead> <tbody> <tr> <td>gtx960</td> <td>16</td> <td>51%</td> <td>60.73%</td> <td>171%</td> <td>167.33%</td> </tr> <tr> <td>gtx960</td> <td>8</td> <td>59%</td> <td>72.03%</td> <td>187%</td> <td>155.25%</td> </tr> <tr> <td>gtx1080</td> <td>16</td> <td>42%</td> <td>41.34%</td> <td>207%</td> <td>137.52%</td> </tr> <tr> <td>rtx2060s</td> <td>16</td> <td>49%</td> <td>57.53%</td> <td>211%</td> <td>149.48%</td> </tr> <tr> <td>rx560</td> <td>16</td> <td>53%</td> <td>56.82%</td> <td>153%</td> <td>115.63%</td> </tr> <tr> <td>rx560</td> <td>8</td> <td>55%</td> <td>54.19%</td> <td>172%</td> <td>122.64%</td> </tr> <tr> <td>intel-hd530</td> <td>8</td> <td></td> <td></td> <td>109%</td> <td>66.12%</td> </tr> </tbody> </table> <p>אומנם זו התחלה אבל כבר התחלה טובה!</p> </div> רשתות נוירונים בקוד פתוח... תמונת מצב http://artyom.cppcms.com/post/328 http://artyom.cppcms.com/post/328 <div style="direction:rtl"> <p>כידוע היום שוק ה־deep learning נשלט באופן <a href="http://artyom.cppcms.com/post/324">כמעט בלעדי</a> ע"י nVidia. אומנם כל תשתיות למידה החישובית הפופולריות כגן TensorFlow, PyTorch, Caffe, MXNet ואחרות משוחררות כקוד פתוח, אבל בליבו של כל אחד מהם, ללא יוצא מן הכלל, רצות ספריות cublas ו־cudnn המאפשרות לנצל את החומרה בצורה מיטבית. כולן כמובן קוד בסגור ומסוגר הרץ על בסיס CUDA. כמובן, גם הוא API פרטי וקנייני של חברת nVidia.</p> <p>אקדים ואומר: אין אני טוען שהסכנה כאן כי החברה "המרושעת" תשתלט על בינה מלאכותית ותקים skynet מתחת לרגליים שלנו. לא, בסה"כ מדובר במימוש פעולות מתמטיות בסיסיות מוגדרות היטב בצורה יעילה להפליא.</p> <h2>אבל אני רוצה קוד פתוח?</h2> <p>אז יש מספר פתרונות וכיוונים:</p> <ol> <li>לאמן הכל ב־CPU בלבד.</li> <li>להשתמש בתשתית ROCm של AMD.</li> <li>להשתמש ב־OpenCL במקום ב־CUDA ואז חוץ מדרייבר של nVidia הכל יהיה פתוח (פחות או יותר)</li> </ol> <p> <a href="/post/328">המשך...</a> </p> </div> על למידה חישובית, תכנה חופשית ומה שביניהם http://artyom.cppcms.com/post/324 http://artyom.cppcms.com/post/324 <div style="direction:rtl"> <p>רשתות נוירונים מהווים היום את שיטת הלמידה החשובה ביותר. הם הביאו לפרצות דרך חשובות. היום כל אדם בעל ידע בתכנות ורקע מתמטי סביר יכול לממש דברים שהיו מדע בדיוני לפני עשור. כוח החישוב העצום של מעבדים גרפיים וזמינות גבוהה של נתונים שינה את פני למידה החישובית. היום אם אתה רוצה להתעסק תחום ראיה ממוחשבת, עיבוד קוד תרגומים וכד' חייב להכיר את השיטות האלה.</p> <p>היום קיימות עשרות תשחתיות (frameworks) לעבודה עם רשתות נוירונים - וכל הפופולריים ביניהם הם תכנה חופשית: tensorflow, pytorch, caffe, keras, mxnet ועוד רבים אחרים הם תכנה חופשית שמופצת תחת רשיונות די מתרניים. חברות ענק שעומדות מאוחרי חלק מהם כמו facebook ו־google דואגים להחזיק את הקוד הפתוח - כי רק כך ניתן לשרוד בעולם הזה בו השיטות והמאמרים שפורסמו לפני שנה כבר לא מספיק עדכניים.</p> <p>אבל, יש פה אבל אחד גדול מאוד. כל התשתיות האלה, דורשות שימוש ב־GPU על מנת לקבל תוצאות בזמן סביר. נקח לדוגמה את הרשת המקורית הידועה בשם alex-net שהייתה אחת פרצות הדרך מהמשעותיות ביותר בתחום הלמידה החישובי בשני עשורים אחרונים. זמן אימון הרשת ב־2012 לקח סדר גודל של שבועיים תוך שימוש בשני כרטיסים גרפיים.</p> <p>כמובן אין כל פסול בשימוש בכרטיסים גרפיים - הם בסה"כ עושים מה שהם יודעים לעשות טוב number-crunching. אבל, היום כמט כל התשתיות מסתמכות של טכנולוגיה אחת וספק אחד - כולם משתמשים ב־cuda וב־nVidia. יתרה מזו חלק מהתשתיות מסתכמות באופן בלעדי על ספריה סוגרה אחת בשם cuDNN שמאפשרת לנצל את כל החישוב של החומרה עד תום. cuDNN ו־cuBLAS הן הספריות שבלעדיהם tensorflow או pytorch פשוט לא יכולים להקיים מבחינת לקוח הקצה.</p> <p>כן, קיימות תשתיות שמאפשרות אימון גם על טכנולוגיה פתוחה. לדוגמה ל־caffe יש ענף opencl העובד על בסיס טכנולוגיות פתוחות ויודע לרוץ גם על כרטיסים של AMD ואפילו של Intel. אבל</p> <ol> <li>פיתוח של caffe די נפסק - ובעולם הדינאמי של היום זה אומר - הפרויקט במצב מוות קליני</li> <li>גם כשאתה משתמש בו אתה מקבל קנס לא קטן מבחינת ביצועיים. זמני הריצה הם איטיים בערך פי שתיים.</li> </ol> <p>בהתחשב בעובדה שחלק מהאימונים יכולים לקחת שעות רבות אפילו ימים זה הופך את הענף של opencl לפחות רלוונטי. הסיבה לאיטיות היא שהמימוש לא נהנה האופטימיזציות מטורפות וכתיבה ב־assembly ש־nVidia הייתה יכולה לעשות ב־cudnn ו־cuBlas.</p> <p>אבל מה עם AMD? האם הם ישנים? כן ולא. AMD דאגו לפתח אלטרנטיבה בשם ROCm. למעשה אם אתה עובד על לינוקס ויש לך כרטיס כמו rx580 או Vega 56 אתה יכול באמץ סביר להריץ את ה־tensorflow ו־pytorch ואפילו caffe על AMD. והיתרון הגדול של ROCm הוא שמדובר בקוד פתוח לחלוטין. החסרון?.. מאיפה להתחיל</p> <ol> <li>ROCm תומך אך ורק בלינוקס אם אתה על Mac או על Windows... לא</li> <li>ספריית MIOpen שלהם שמהווה מאין תחליף ל־cudnn, אפילו שתומכת ב־OpenCL עובדת אך ורק על דיריבר rocm של AMD. משמעות - אומנם זה קוד פתוח אבל זה vendor-lock-in לא פחות מ־cudnn של nvidia</li> <li>ROCm לא תומך עדיין ברטיסים הגרפיים העדכניים ביותר מבוססי rdna כמו Rx 5700XT וחבריו. עברה שנה מאז שהכטריסים האלה הושקו אבל עדיין לא ניתן להשתמש בהם לטובת למידה חישובית.</li> <li>הוא גם לא נותן מענה ל־APUs. הכרטיסים הגרפיים המובנים שבאים במעבדים כמו Razen 3400G - לא יעבדו עם tensorflow או pytorch. ויש לציין של־Vega 11 שבא עם 3400G יש יותר כוח החישוב מ־GTX 580 ש־alex-net המקורי אומן עליו.</li> </ol> <p>למעשה נראה כי AMD עשתה הכל כדי למנוע ממישו אפילו להסתכל בכיוון שלהם לטובת deep-learnים.</p> <p>מה עם פתרונות עבור intel? הרי גם להם יש GPU? מעבר לעובד שביצועי Intel GPU הם בדיחה, גם intel דאגה לכתוב ספריית deep-learning משלה שלא עובדת עם שום כרטיס גרפי אחר.</p> <h2>שורה תחתונה</h2> <p>למרות שמבחוץ נראה שכל נושא למידה חישובית על רשתות נוירונים מתבסס על תכנה חופשית, במציאות יש רק דרך אחת לעבוד - לעבוד עם הקוד הסגור של ספק אחד. ללא שילוב של nVidia/cuda/cudnn התחזיות של Deep-Learning די עגומות</p> </div> בינה מלאכותית על ZX Spectrum http://artyom.cppcms.com/post/322 http://artyom.cppcms.com/post/322 <div style="direction:rtl"> <p>המחשב הראשון שלי היה <a href="https://he.wikipedia.org/wiki/%D7%A1%D7%A4%D7%A7%D7%98%D7%A8%D7%95%D7%9D_%D7%A1%D7%99%D7%A0%D7%A7%D7%9C%D7%99%D7%99%D7%A8">ZX Spectrum</a>. למדתי עליו לתכנת, למדתי לכתוב קוד אסמבלי, וגם את ההבנה איך מעבדים בנויים רכשתי שם.</p> <p>הוא היה הכלי ששימש אותי בכתיבת סימולציות פיזיקאליות בזמן שלמדתי בבית ספר עם דגש בתחום פיזיקה ומתמטיקה. אפילו אחי הגדול שלמד באוניברסיטה כתב עליו חישובים מתמטיים מסובכים לטובת הלימודים. זו הייתה האהבה הדיגיטלית הראשונה שלי.</p> <p>היום אני עוסק בתחום בבינה מלאכותית ומשתמש בכרטיסים גרפיים חזקים ביותר שהביצועים שלהם נמדדים ב־Terra FLOPS. אבל לאחרונה נתקעתי בסימולטור של ZX Spectrum ועלה במוחי רעיון. האם אפשר לקחת את המשימות שאני עושה היום ולעשות אותה על המחשב של אז?</p> <p>אז לקחתי את ה-Hello World של למידה חישובית <a href="https://en.wikipedia.org/wiki/MNIST_database">זיהוי ספרות בכתב היד</a> והחלטתי לממש את זה ב-ZX Spectrum.</p> <p>להלן התוצאות:</p> <p><a href="https://github.com/artyom-beilis/zx_spectrum_deep_learning">https://github.com/artyom-beilis/zx_spectrum_deep_learning</a></p> <p><img src="https://user-images.githubusercontent.com/14816918/71548763-9db92080-29bb-11ea-82a9-34cd1a510c26.png" alt="mnist2" /></p> <p>ניתן למצוא מאמר מלא באנגלית והסברים מלאים על התהליך כאן:</p> <p><a href="http://blog.cppcms.com/post/125">http://blog.cppcms.com/post/125</a></p> </div>