منذ وقت ليس ببعيد تحدثنا عن خدمة Selectel الجديدة -
الحوسبة السحابية عالية الأداء على مسرعات FPGA . في مقالة جديدة حول هذا الموضوع ، نعتبر مثالًا لبرمجة FPGA لإنشاء مجموعة Mandelbrot ، وهي خوارزمية رياضية معروفة لتصور صور كسورية. استخدمت المقالة مادة من موقع
مشروع أويلر .
بدلا من المقدمة
أولاً ، بعض المصطلحات. نظام كمبيوتر مع مسرع FPGA - كقاعدة ، هذا محول PCIe مع شريحة FPGA كجزء من خادم x64. يأخذ المسرع مهمة منفصلة كثيفة الموارد يمكن أن تشارك فيها الحوسبة المتوازية وينفذها العديد من الطلبات بحجم أسرع من معالج x64 ، وتفريغها وزيادة أداء نظام الحوسبة بأكمله. على سبيل المثال ، يمكن إجراء دورة حسابية تحتوي على 100 ألف تكرار في FPGA بتمريرة واحدة فقط ، بدلاً من تنفيذ 100 ألف مرة على التوالي على معالج x64 كلاسيكي. العناصر المنطقية وموارد الأجهزة وروابط الاتصال وشرائح FPGA مبرمجة من قبل المستخدم مباشرة للمهمة نفسها ، والتي تتيح لك تنفيذ المهمة كتطبيق خوارزمية في السيليكون - خوارزمية في السيليكون وبالتالي تحقيق أداء عالي ، ومع استهلاك طاقة متواضع للغاية.
اليوم ، عتبة الدخول إلى تقنية FPGA متاحة تمامًا حتى للشركات الناشئة - يمكن استئجار خادم مزود بمسرع FPGA وجميع البرامج الضرورية (SDK) في سحابة Selectel بسعر معقول (ما يسمى بـ "cloud FPGA") ، ويؤدي دعم معيار Open CL في FPGA إلى أن المبرمج الذي يعرف كيفية العمل مع C قادر على إعداد وتشغيل برنامج على FPGA.
التطلع إلى المستقبل: جرب FPGA في العمل
المثال البرمجي الموصوف أدناه لبناء مجموعة ماندلبروت تم تنفيذه بالفعل على خادم اختبار
في معمل سيليكتل ، حيث يمكن لأي شخص تقييم أدائه (سيكون التسجيل مطلوبًا).
يتم توفير المشروع في التعليمات البرمجية وإعداده لتجميع. توفر Selectel وصولاً بعيدًا إلى خادم مزود بمسرع Intel Arria 10 FPGA. على جانب الخادم ، يتم نشر أدوات SDK و BSP لتطوير OpenCL وتصحيحه وتجميعه ، ويتم نشر كود Visual Studio لإعداد التطبيقات المضيفة (تطبيقات التحكم للمعالج المركزي للخادم).
لاحظ أن المثال نفسه ليس له أي قيمة تطبيقية ؛ تم اختياره لأسباب شرح طرق التسريع باستخدام مبادئ التوازي. في هذا المثال ، يتعرف القارئ على مسار تصميم تطبيق في نظام حوسبة غير متجانس باستخدام FPGA - يمكن استخدام هذا المسار لاحقًا لتطوير تطبيقاتك باستخدام حوسبة متوازية.
تحديث : في ربيع عام 2018 ، قدمت Intel المعالج الهجين عالي الأداء Xeon Gold 6138P مع شريحة Arria 10 FPGA مدمجة. بحلول نهاية عام 2018 ، من المتوقع أن تصبح المعالجات التسلسلية من هذا النوع متاحة للعملاء من خلال شركاء Intel. نحن في Selectel نتطلع إلى هذه الشريحة ، ونأمل أن نكون الأوائل في روسيا الذين يوفرون لعملائنا الفرصة لاختبار هذا المنتج الجديد الفريد.
حول معيار OpenCL لبرمجة FPGA
تم تطوير معيار OpenCL من قبل مجموعة Khronos Group ، كبرى الشركات المصنعة للرقائق والبرامج في العالم والتي تضم Intel و AMD و Apple و ARM و Nvidia و Sony Computer Entertainment وغيرها. وقد تم تصميمه لكتابة التطبيقات التي تستخدم الحوسبة المتوازية على أنواع مختلفة من المعالجات ، بما في ذلك FPGA. يتضمن معيار OpenCL لغة البرمجة C بناءً على إصدار لغة C99 (أحدث إصدار من C99 هو ISO / IEC 9899: 1999 / Cor 3: 2007 من 2007-11-15) وبيئة برمجة التطبيقات.
تعتمد شعبية استخدام OpenCL للحوسبة عالية الأداء على حقيقة أنه معيار مفتوح ، ولا يتطلب استخدامه ترخيصًا. علاوة على ذلك ، لا يقصر OpenCL نطاق الأجهزة المدعومة على أي علامة تجارية معينة ، مما يسمح باستخدام الأجهزة من الشركات المصنعة المختلفة على نفس النظام الأساسي للبرامج.
بالإضافة إلى OpenCL: مقدمة إلى OpenCL عن هبر .
القليل من التاريخ - كان مسار تصميم FPGA الذي كان موجودًا قبل معيار OpenCL محددًا للغاية ومستهلكًا للوقت ، بينما كان من حيث التعقيد متفوقًا على تصميم الرقائق المخصصة (ASIC ، الدائرة المتكاملة الخاصة بالتطبيقات ، "الدائرة المتكاملة ذات الأغراض الخاصة"). كان مطلوبًا فهمًا شاملاً لبنية أجهزة FPGA ، وكان يجب إجراء تكوينها بلغة وصف الأجهزة منخفضة المستوى (HDL). كان امتلاك هذا التصميم والتحقق مسارًا ولا يزال فنًا ، نظرًا للتعقيد الشديد ، متاحًا لدائرة محدودة من المطورين.
لقد أدى ظهور مجموعة أدوات دعم OpenCL من Intel لـ FPGAs جزئيًا إلى معالجة مسألة إمكانية الوصول إلى برمجة FPGA لمطوري البرامج. يختار المبرمج بشكل مستقل ذلك الجزء من الخوارزمية المناسب للمعالجة المتوازية ويصفه في C ، ثم يقوم مترجم Intel OpenCL لـ FPGA بإنشاء ملف تكوين ثنائي لتشغيل هذا الجزء من الخوارزمية على المسرّع.
باستخدام بيئة Visual Studio المعتادة أو مترجم gcc القياسي ، يتم تحضير تطبيق مضيف (تطبيق من نوع .exe ، يتم تنفيذه على معالج x64 الرئيسي) ، بينما يتم تضمين جميع مكتبات الدعم الضرورية في SDK. عندما يتم تشغيل التطبيق المضيف ، يتم تحميل البرامج الثابتة FPGA ، سيتم تحميل البيانات في قلب الشريحة وستبدأ المعالجة وفقًا للخوارزمية المتصورة.
FPGA (FPGA) عبارة عن بنية أجهزة متوازية هائلة قابلة للبرمجة من قِبل المستخدم مع ملايين العناصر المنطقية وآلاف وحدات إشارة DSP وعشرات الميغابايت من ذاكرة التخزين المؤقت للحسابات المدمجة ، دون الوصول إلى وحدات الذاكرة الرئيسية للخادم. تتيح لك واجهات الإدخال / الإخراج السريعة (10GE و 40GE و 100GE و PCIe Gen 3 وما إلى ذلك) تبادل البيانات بفعالية مع المعالج الرئيسي للخادم.
معيار OpenCL هو بيئة لتنفيذ برامج غير متجانسة. تتكون البيئة من جزأين منفصلين:
- برنامج مضيف - تطبيق يعمل على المعالج المركزي الرئيسي للخادم ، مكتوب بلغة C / C ++ وباستخدام مجموعة وظائف OpenCL API. ينظم الخادم المضيف العملية الكاملة للحوسبة ، وتزويد المصدر وتلقي بيانات الإخراج ، ويتفاعل مع جميع أنظمة الخادم باستخدام مسرع FPGA.
- برنامج Accelerator - برنامج مكتوب بلغة OpenCL C (لغة C مع عدد من القيود) ، تم تجميعه للتشغيل على شريحة FPGA.
الخادم النموذجي للحوسبة المتوازية هو جهاز كمبيوتر يستند إلى x64 (لتشغيل التطبيقات المضيفة) ، والذي يتضمن مسرع FPGA للأجهزة ، والذي يتم توصيله غالبًا عبر ناقل PCI-Express. بالمناسبة ، يتم عرض مثل هذا النظام في مختبر Selectel.
يتكون تسلسل البرمجة والتجميع لمسرع FPGA من مرحلتين. يتم تجميع رمز التطبيق المضيف بواسطة مترجم قياسي (Visual C ++ ، GCC) للحصول على ملف قابل للتنفيذ في نظام تشغيل الخادم (على سبيل المثال ، * .exe). يتم تحضير الكود المصدري لمسرع FPGA (kernel، kernel) بواسطة برنامج التحويل البرمجي AOC كجزء من SDK ، مع استلام ملف ثنائي (* .aocx). هذا الملف مخصص فقط لبرمجة المعجل.
التين. بنية بيئة تجميع البرمجيات OpenCL
ضع في اعتبارك بعض الأمثلة البرمجية لحساب ناقلات كبيرة بطريقتين
(
ملاحظة: لا تقم بإطلاق النار على عازف البيانو - فيما يلي يتم استخدام الرمز من موقع مشروع أويلر ):
void inc (float *a, float c, int N) { for (int i = 0; i<N; i++) a[i] = a[i] + c; } void main() { ... inc(a,c,N); ... }
_kernel void inc (_global float *a, float c) { int i = get_global_id(0); a[i] = a[i] + c; } void main() { ... clEnqueueNDRangeKernel(...,&N,...) ... }
الشفرة في البداية هي مثال لكيفية ظهور التنفيذ المفرد في لغة C باستخدام طريقة الحساب المتسلسل للعناصر العددية.
الإصدار الثاني من الكود هو تطبيق محتمل للخوارزمية على OpenCL في شكل دالة محسوبة على مسرّع FPGA. لا يوجد حلقة ، ويتم الحساب في تكرار واحد للحلقة. يحدث حساب صفيف متجه عند تنفيذ نسخ N من هذه الوظيفة. كل نسخة لها فهرس خاص بها ، يتم استبداله في المكرر في حلقة ، ويتم تعيين عدد مرات إعادة المحاولة من المضيف عند تنفيذ التعليمات البرمجية. يتم توفير إجراء المكرر بواسطة دالة get_global_id () ، والتي تعمل مع فهرس داخل مؤشر 0 ≤ <N.
نصل إلى النقطة: بناء فراكتل
مجموعة Mandelbrot عبارة عن مصفوفة من النقاط "c" على المستوى المعقد والتي تحدد لها علاقة التكرار Zn + 1 = Zn² + c لـ Z0 = 0 تسلسل محدد.
نحدد Zn = Zn + IYn ، وكذلك c = p + iq.
يتم حساب التسلسل التالي لكل نقطة:
Xn + 1 = Xn² + Yn² + ص
Yn + 1 = 2XnYn + ف
يتم حساب انتماء نقطة إلى المجموعة عند كل تكرار كمعادلة
Xn² + Yn² <4.
لعرض مجموعة ماندلبروت على الشاشة ، نحدد قاعدة:
- إذا استمر عدم المساواة في أي تكرار ، فستدخل النقطة إلى المجموعة وستظهر باللون الأسود.
- إذا لم يثبت التفاوت ، بدءًا من قيمة تكرار معينة n = N ، فسيتم تحديد اللون من خلال عدد التكرارات N.
ستكون العملية الحسابية على المضيف كما يلي:
- يتم تعيين حساب عدد التكرارات لكل نقطة داخل نافذة البكسل لوظيفة mandel_pixel ().
- سيتم توفير التعداد التسلسلي لنقاط الصورة بواسطة وظيفة softwareCalculateFrame (). تحدد المعلمات الفاصل الزمني الحقيقي للنقاط المحسوبة ، والخطوة الحقيقية للخوارزمية ، ومؤشرًا لمخزن الألوان لحجم الصورة (theWidth * theHeight).
- يتم تعديل لون النقطة بواسطة SoftColorTable.
دعنا ننتقل إلى الرمز:
inline unsigned int mandel_pixel( double x0, double y0, unsigned int maxIterations ) { // variables for the calculation double x = 0.0; double y = 0.0; double xSqr = 0.0; double ySqr = 0.0; unsigned int iterations = 0; // perform up to the maximum number of iterations to solve // the current point in the image while ( xSqr + ySqr < 4.0 &&iterations < maxIterations ) { // perform the current iteration xSqr = x*x; ySqr = y*y; y = 2*x*y + y0; x = xSqr - ySqr + x0; // increment iteration count iterations++; } // return the iteration count return iterations; }
int softwareCalculateFrame( double aStartX, double aStartY, double aScale, unsigned int* aFrameBuffer ) { // temporary pointer and index variables unsigned int * fb_ptr = aFrameBuffer; unsigned int j, k, pixel; // window position variables double x = aStartX; double y = aStartY; double cur_x, cur_y; double cur_step_size = aScale; // for each pixel in the y dimension window for ( j = 0, cur_y = y; j < theHeight; j++, cur_y -= cur_step_size ) { // for each pixel in the x dimension of the window for ( cur_x = x, k = 0; k< theWidth; k++, cur_x += cur_step_size ) { // set the value of the pixel in the window pixel = mandel_pixel(cur_x, cur_y, theSoftColorTableSize); if ( pixel == theSoftColorTableSize ) *fb_ptr++ = 0x0; else *fb_ptr++ = theSoftColorTable[pixel]; } } return 0; }
يتم حساب كل بكسل بشكل مستقل عن الآخر ، وبالتالي يمكن موازاة هذه العملية. عند تنفيذ الخوارزمية لتسريع FPGA ، يتم إنشاء تعليمات SIMD لحساب الرقم لكل بكسل تكرار (تحديد رمز اللون من اللوحة). يتم تأطير تنفيذ حلقتين متداخلين على المخزن المؤقت للصور من خلال OpenCL عن طريق تشغيل العملية (theWidth * theHeight).
تُسمى مثيلات kernel في القائمة أدناه عنصر العمل ، وتسمى مجموعة جميع المثيلات مساحة الفهرس. تتضمن ميزات وظيفة الجهاز ما يلي:
- يبدأ تعريف الوظيفة بالكلمة الرئيسية __kernel.
- نوع وظيفة الأجهزة - نوع القيمة المرتجعة باطل دائمًا.
- يتم إرجاع القيم من خلال المخازن المؤقتة التي تم تمريرها كمعلمات.
- تحدد المعلمات الثلاثة الأولى شبكة المواد ، التي تتوافق عُقدها مع وحدات البكسل لصورة الإخراج.
- تحدد المعلمة الرابعة عدد التكرارات ، مما يمنع التكرار للنقاط التي تنتمي إلى مجموعة Mandelbrot.
- المعلمة الخامسة هي مؤشر إلى المخزن المؤقت للألوان الناتج.
- تشير الكلمة الأساسية __global إلى نوع الذاكرة التي سيتم من خلالها إرسال المخزن المؤقت: هذه هي ذاكرة DDR (QDR) العامة على المسرّع نفسه.
- تعطي الكلمة الأساسية المقيدة للمحسن حظراً على استخدام مراجع المخزن المؤقت غير المباشر.
- في المعلمة السادسة ، يتم تمرير مؤشر لوح الألوان.
- تعمل الكلمة الرئيسية __constant على تحسين الوصول إلى المخزن المؤقت عن طريق إنشاء ذاكرة تخزين مؤقت بسمة للقراءة فقط.
وصف الوظيفة في القائمة قريب من تنفيذ معالج x64. هنا ، يتم تعريف مثيل kernel الحالي من خلال دالة get_global_id ، حيث يتم تمرير رقم البعد (0 ، 1) كمعلمة.
لتحسين التحسين ، تم تقديم إشارة صريحة لبدء الدورة. في حالة عدم وجود معلومات حول عدد التكرارات في وقت التجميع ، يتم تحديد عدد خطوات الحلقة بشكل صريح ، حيث سيتم إنشاء مجموعات الأجهزة الخاصة بها. مع هذا النوع من الترميز ، يجب على المرء أن "ينظر إلى الوراء" في قدرة شريحة معينة مثبتة على المسرّع ، وذلك بسبب استهلاك موارد FPGA لعدد أكبر من الدورات.
//////////////////////////////////////////////////////////////////// // mandelbrot_kernel.cl : Hardware implementation of the mandelbrot algorithm //////////////////////////////////////////////////////////////////// // Amount of loop unrolling. #ifndef UNROLL #define UNROLL 20 #endif // Define the color black as 0 #define BLACK 0x00000000 __kernel void hw_mandelbrot_frame ( const double x0, const double y0, const double stepSize, const unsigned int maxIterations, __global unsigned int *restrict framebuffer, __constant const unsigned int *restrict colorLUT, const unsigned int windowWidth) { // Work-item position const size_t windowPosX = get_global_id(0); const size_t windowPosY = get_global_id(1); const double stepPosX = x0 + (windowPosX * stepSize); const double stepPosY = y0 - (windowPosY * stepSize); // Variables for the calculation double x = 0.0; double y = 0.0; double xSqr = 0.0; double ySqr = 0.0;</code> <code>unsigned
يجب تثبيت حزمة أدوات Intel FPGA SDK لحزمة OpenCL على المضيف قبل تجميع تطبيق الأجهزة للخوارزمية. من بين البرامج المثبتة مسبقًا ، يجب عليك تضمين BSP (حزمة دعم اللوحة) من الشركة المصنعة للوحة التسريع المحددة. في المثال ، تم تثبيت Intel Quartus Prime Pro 16.1 مع دعم OpenCL و BSP لمسرع Euler Thread (Intel Arria 10).
فيما يلي تكوين المسارات ومتغيرات البيئة. يحتوي المتغير ALTERAOCLSDKROOT على المسار إلى Intel FPGA SDK ، ويحتوي المتغير AOCL_BOARD_PACKAGE_ROOT على مُسرع BSP.
set ALTERAOCLSDKROOT=C:\intelFPGA_pro\16.1\hld set AOCL_BOARD_PACKAGE_ROOT=C:\intelFPGA_pro\16.1\hld\board\euler_thread set path=%path%;C:\intelFPGA_pro\16.1\hld\bin set path=%path%;C:\intelFPGA_pro\16.1\quartus\bin64 set path=%path%;C:\intelFPGA_pro\16.1\hld\board\a10_ref\windows64\bin set path=%path%;C:\intelFPGA_pro\16.1\hld\host\windows64\bin set path=%path%;C:\intelFPGA_pro\16.1\qsys\bin set path=%path%;C:\Program Files (x86)\GnuWin32\bin\
للترجمة ، يتم استخدام برنامج التحويل البرمجي aoc من SDK.
aoc mandelbrot_kernel.cl -o mandelbrot_kernel.aocx --board thread -v -v --report
نقوم بفك تشفير: mandelbrot_kernel.cl - الملف الذي يحتوي على النص المصدر ، mandelbrot_kernel.aocx - ملف كائن الإخراج لبرمجة FPGA ، خيط - اسم المسرع من حزمة BSP. يعرض رمز التبديل --report تقرير استخدام مورد FPGA. يعرض رمز التبديل –v معلومات التشخيص أثناء التجميع. تقرير استهلاك الموارد للنواة كالتالي:
+ ------------------------------------------------- ------------------- +
؛ ملخص استخدام الموارد التقديرية ؛
+ ---------------------------------------- + -------- ------------------- +
؛ الموارد + الاستخدام ؛
+ ---------------------------------------- + -------- ------------------- +
؛ استخدام المنطق 49٪
؛ ALUTs ؛ 26٪
؛ سجلات المنطق المخصصة ؛ 25٪ ؛
؛ كتل الذاكرة 21٪
؛ كتل DSP ؛ 16٪ ؛
+ ---------------------------------------- + -------- ------------------- ؛
لتجميع التطبيق المضيف ، استخدم المثال حزمة Microsoft Visual Studio 2010 Express مع تثبيت Microsoft SDK 7.1. في إعدادات المشروع ، يتم تحديد تكوين x64. بعد ذلك ، قم بتوصيل المجلد لملفات الرؤوس الخارجية وحدد المسار إلى مكتبات Intel FPGA SDK الإضافية في إعدادات الرابط.
الدلائل الإضافية لتضمين الملفات = $ (ALTERAOCLSDKROOT) \ host \ include؛
أدلة مكتبة إضافية = $ (AOCL_BOARD_PACKAGE_ROOT) \ windows64 \ lib؛
$(ALTERAOCLSDKROOT)\host\windows64\lib;
ستكون خطة العمل العامة لإطلاق النواة على المسرِّع كما يلي:
- احصل على قائمة المنصات
- احصل على قائمة بالأجهزة
- إنشاء سياق ؛
- تحميل النواة في الجهاز ؛
- إرسال مخازن الإدخال إلى الجهاز ؛
- تشغيل النواة للتنفيذ ؛
- قراءة المخزن المؤقت للإخراج من الجهاز ؛
- سياق حر.
ضع في اعتبارك بعض النقاط المتعلقة مباشرة بإطلاق النواة. لذلك ، تم تصميم نواة واحدة لمعالجة بكسل واحد من الصورة. وبالتالي ، تحتاج إلى تشغيل مثيلات N kernel ، حيث N هو العدد الإجمالي لوحدات البكسل في الصورة.
أدناه ، نلاحظ الحالة عندما يكون هناك العديد من لوحات التسريع في الخادم ، ثم يمكن توزيع المهمة بينهما. في كل مسرعات ، تحتاج إلى تحميل النواة (ملف mandelbrot_kernel.aocx). لنفترض أن عدد المسرعات هو numDevices ، وأن خطوط الصورة مقسمة بين جميع المسرعات:
#define MAXDEV 10 static cl_context theContext; static cl_program theProgram; static cl_kernel theKernels[MAXDEV]; //.. // Create the program object theProgram = createProgramFromBinary( theContext, "mandelbrot_kernel.aocx", theDevices, numDevices); // Create the kernels for ( unsigned i = 0; i < numDevices; ++i ) theKernels[i] = clCreateKernel( theProgram, "hw_mandelbrot_frame", &theStatus ); // Create output pixel buffers for every kernel for( unsigned i = 0; i < numDevices; ++i ) thePixelData[i] = clCreateBuffer(theContext, CL_MEM_WRITE_ONLY, thePixelDataWidth*rowsPerDevice[i]*sizeof(unsigned int), NULL, &theStatus); // Preparing and writing palette buffer to every device theHardColorTable = clCreateBuffer(theContext, CL_MEM_READ_ONLY, aColorTableSize*sizeof(unsigned int), NULL, &theStatus); for( unsigned i = 0; i < numDevices; i++ ) theStatus = clEnqueueWriteBuffer(theQueues[i], theHardColorTable, CL_TRUE, 0, aColorTableSize*sizeof(unsigned int), aColorTable, 0, NULL, NULL); // Preparing kernels and run unsigned rowOffset = 0; for ( unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++] ) { // Create ND range size size_t globalSize[2] = { thePixelDataWidth, rowsPerDevice[i] }; // Set the arguments unsigned argi = 0; theStatus = clSetKernelArg (theKernels[i], argi++, sizeof(cl_double), (void*) &aStartX ); const double offsetedStartY = aStartY - rowOffset * aScale; theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_double), (void*)&offsetedStartY); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_double), (void*)&aScale); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theHardColorTableSize); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&thePixelData[i]); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&theHardColorTable); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theWidth); // Launch kernel theStatus = clEnqueueNDRangeKernel(theQueues[i], theKernels[i], 2, NULL, globalSize, NULL, 0, NULL, NULL); } rowOffset = 0; for( unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++] ) { // Read the output theStatus = clEnqueueReadBuffer(theQueues[i], thePixelData[i], CL_TRUE, 0, thePixelDataWidth*rowsPerDevice[i]*sizeof(unsigned int), &aFrameBuffer[rowOffset * theWidth], 0, NULL, NULL); } / / . .
- تنشئ الدالة createProgramFromBinary كائن برنامج OpenCL من ملف كائن.
- بعد ذلك ، يتم إنشاء نواة لكل جهاز استنادًا إلى كائن البرنامج.
- يتم إنشاء مخازن PixelData لتلقي الإخراج من كل قلب.
- يتم إنشاء مخزن مؤقت لتخزين لوحة الألوان وتحميله في كل مسرعات.
- بعد ذلك ، لكل جهاز ، يتم تعيين ربط معلمات التطبيق المحلي ومعلمات kernel باستخدام وظيفة clSetKernelArg.
- يتم تحديد المعلمات بالأرقام التسلسلية في إعلان دالة kernel ، بدءًا من الصفر.
النقطة المهمة التالية هي تحديد حجم المهمة بناءً على مساحة الفهرس وفقًا لصفيف globalSize. يمكن أن يكون هذا الصفيف أحاديًا أو ثنائيًا أو ثلاثي الأبعاد. لكل بعد ، يتم إعطاء البعد كعدد صحيح. سيحدد بُعد المساحة ترتيب فهرس عنصر العمل في النواة.
في المثال ، يتم تحديد مسافة ثنائية الأبعاد لكل قلب ، حيث يكون أحد المحاور هو عناصر صف البكسل ، والثاني هو مجموعة خطوط الصورة التي تتم معالجتها على هذا الجهاز. في كود النواة ، يتم الحصول على رقم البكسل في الخط عن طريق استدعاء get_global_id (0) ، ورقم السطر هو get_global_id (1). يتم تمرير متغير globalSize إلى الدالة clEnqueueNDRangeKernel لبدء العدد المطلوب من مثيلات kernel المطلوب تنفيذها.
عند الانتهاء من تنفيذ النوى ، تتم قراءة مخازن البكسل من الجهاز إلى المصفوفات المحلية. لنقم بتقييم الأداء بعدد الإطارات في الثانية - تكون النتيجة مرئية على العرض التوضيحي الذي تم إجراؤه في مؤتمر SelectelTechDay ( انظر بداية المقالة ).
الخلاصة
أدت برمجة مسرعات FPGA بلغة عالية المستوى دون شك إلى خفض عتبة الوصول إلى هذه التقنية للمطورين بترتيب من الحجم. على سبيل المثال ، بالنسبة لأولئك الذين يتقنون مجموعة الأدوات هذه ، يوجد حتى تطبيق FPGA لمثال "Hello World" الشهير.
ولكن ليس بهذه البساطة. الكتابة - وخاصة - تصحيح الخوارزمية التي تعمل بشكل واضح لمشكلة تطبيقية حقيقية لا تزال تتطلب احترافية عالية. قيود أخرى هي أن كل شريحة FPGA يمكن أن تؤدي مهمة حسابية واحدة فقط داخل التطبيق. لمهمة أخرى ، يجب إعادة برمجته مرة أخرى.
بالمناسبة ، يتيح لك نموذج استخدام النظام الأساسي وجود أكثر من مسرع FPGA على المضيف ، على الرغم من أن هذا حل مكلف إلى حد ما.
يدير المضيف (التطبيق المضيف) عملية إنشاء السياق (هيكل البيانات للمسرع) وقائمة انتظار الأوامر. على سبيل المثال تطبيق مضيف واحد ، حيث توجد مهام فرعية مختلفة للحوسبة المتوازية على FPGA ، يمكن تحميلها على مسرعات مختلفة:
KERNEL1 => المسرع أ
KERNEL2 => المسرع ب
ومع ذلك ، فإن الجهود المبذولة لإتقان مسرعات FPGA تستحق العناء - في العديد من مجالات التطبيق ، أصبحت هذه التكنولوجيا لا غنى عنها: الاتصالات والتكنولوجيا الحيوية ومعالجة البيانات الكبيرة والتعرف على الأنماط والإشارة ومعالجة الصور في الرياضيات الحاسوبية ونمذجة المجال المادي.
معلومات إضافية للمقال:
www.altera.com هو مورد Intel FPGA الأساسي.
www.eulerproject.com هو الموقع الرسمي لمشروع أويلر.
Altera + OpenCL: نقوم بالبرمجة تحت FPGA دون معرفة VHDL / Verilog - مقالة عن هبر.