Geekly articles weekly

Beispiel für die Programmierung eines FPGA-Beschleunigers



Vor nicht allzu langer Zeit haben wir über den neuen Selectel-Dienst gesprochen - Cloud-Hochleistungsrechnen auf FPGA-Beschleunigern . In einem neuen Artikel zu diesem Thema betrachten wir ein Beispiel für die FPGA-Programmierung zur Erstellung eines Mandelbrot-Sets, eines bekannten mathematischen Algorithmus zur Visualisierung fraktaler Bilder. Der Artikel verwendete Material von der Site Euler Project .



Anstelle des Vorworts


Zunächst einige Begriffe. Ein Computersystem mit einem FPGA-Beschleuniger - in der Regel ist dies ein PCIe-Adapter mit einem FPGA-Chip als Teil des x64-Servers. Der Beschleuniger übernimmt eine separate ressourcenintensive Aufgabe, an der paralleles Rechnen beteiligt sein kann, und führt sie viele Größenordnungen schneller als der x64-Prozessor aus, entlädt sie und erhöht die Leistung des gesamten Computersystems. Beispielsweise kann ein Berechnungszyklus mit 100.000 Wiederholungen auf einem FPGA in nur einem Durchgang ausgeführt werden, anstatt auf einem klassischen x64-Prozessor 100.000 Mal nacheinander ausgeführt zu werden. Logische Elemente, Hardwareressourcen, Kommunikationsverbindungen und FPGA-Chips werden vom Benutzer direkt für die Aufgabe selbst programmiert. Auf diese Weise können Sie die Aufgabe als Implementierung eines Algorithmus in Silizium implementieren - Algorithmus in Silizium - und dadurch eine hohe Leistung und einen sehr geringen Stromverbrauch erzielen.

Die Schwelle für den Einstieg in die FPGA-Technologie ist heute auch für Startups durchaus zugänglich - ein Server mit einem FPGA-Beschleuniger und der gesamten erforderlichen Software (SDK) kann in der Selectel-Cloud zu einem angemessenen Preis gemietet werden (das sogenannte „Cloud-FPGA“), und die Unterstützung des Open CL-Standards im FPGA führt dazu dass ein Programmierer, der weiß, wie man mit C arbeitet, ein Programm auf FPGA vorbereiten und ausführen kann.

Mit Blick auf die Zukunft: Probieren Sie FPGA bei der Arbeit aus


Das unten beschriebene Programmierbeispiel zum Erstellen eines Mandelbrot-Sets wurde bereits auf einem Testserver im Selectel Lab implementiert , auf dem jeder seine Leistung bewerten kann (eine Registrierung ist erforderlich).


Das Projekt wird im Code bereitgestellt und für die Kompilierung vorbereitet. Selectel bietet Fernzugriff auf einen Server mit Intel Arria 10 FPGA-Beschleuniger. Auf der Serverseite werden SDK- und BSP-Tools zum Entwickeln, Debuggen und Kompilieren von OpenCL, Visual Studio-Code zum Vorbereiten von Hostanwendungen (Steueranwendungen für den Zentralprozessor des Servers) bereitgestellt.
Beachten Sie, dass das Beispiel selbst keinen angewendeten Wert hat und aus Gründen der Demonstration der Beschleunigungsmethoden unter Verwendung der Prinzipien der Parallelität ausgewählt wurde. In diesem Beispiel lernt der Leser den Weg zum Entwerfen einer Anwendung in einem heterogenen Computersystem mit FPGA kennen. Später kann dieser Weg verwendet werden, um Ihre eigenen Anwendungen mit parallelem Rechnen zu entwickeln.
UPDATE : Im Frühjahr 2018 stellte Intel den Hochleistungs-Hybridprozessor Xeon Gold 6138P mit integriertem Arria 10 FPGA-Chip vor. Bis Ende 2018 sollen serielle Prozessoren dieses Typs über Intel-Partner für Kunden verfügbar sein. Wir bei Selectel freuen uns auf diesen Chip und hoffen, als erster in Russland unseren Kunden die Möglichkeit zu geben, dieses einzigartige neue Produkt zu testen.

Informationen zum OpenCL-Standard für die FPGA-Programmierung


Der OpenCL-Standard wurde von der Khronos Group entwickelt, den weltweit führenden Chip- und Softwareherstellern, darunter Intel, AMD, Apple, ARM, Nvidia, Sony Computer Entertainment usw. Er wurde zum Schreiben von Anwendungen entwickelt, die Parallel Computing auf verschiedenen Prozessortypen, einschließlich FPGA, verwenden. Der OpenCL-Standard enthält die Programmiersprache C basierend auf der Sprachversion C99 (die neueste Version von C99 ist ISO / IEC 9899: 1999 / Cor 3: 2007 vom 15.11.2007) und eine Anwendungsprogrammierumgebung.

Die Popularität der Verwendung von OpenCL für Hochleistungsrechner beruht auf der Tatsache, dass es sich um einen offenen Standard handelt und für dessen Verwendung keine Lizenz erforderlich ist. Darüber hinaus beschränkt OpenCL die Palette der unterstützten Geräte nicht auf eine bestimmte Marke, sodass Hardware verschiedener Hersteller auf derselben Softwareplattform verwendet werden kann.

Zusätzlich zu OpenCL: Einführung in OpenCL auf Habr .


Ein bisschen Geschichte - die FPGA-Entwurfsroute, die vor dem OpenCL-Standard existierte, war äußerst spezifisch und zeitaufwändig, während sie hinsichtlich der Komplexität dem kundenspezifischen Chipdesign (ASIC, anwendungsspezifische integrierte Schaltung, "spezielle integrierte Schaltung") sogar überlegen war. Ein gründliches Verständnis der FPGA-Hardwarestruktur war erforderlich, deren Konfiguration in einer Low-Level-Hardwarebeschreibungssprache (HDL) durchgeführt werden musste. Der Besitz dieser Entwurfs- und Verifizierungsroute war und ist eine Kunst, die aufgrund der extremen Komplexität einem begrenzten Kreis von Entwicklern zur Verfügung steht.

Das Aufkommen des OpenCL-Support-Toolkits von Intel für FPGAs hat sich teilweise mit dem Problem der Zugänglichkeit der FPGA-Programmierung für Softwareentwickler befasst. Der Programmierer wählt unabhängig den Teil seines Algorithmus aus, der für die Parallelverarbeitung geeignet ist, und beschreibt ihn in C. Anschließend erstellt der Intel OpenCL-Compiler für FPGA eine binäre Konfigurationsdatei, um dieses Fragment des Algorithmus auf dem Beschleuniger auszuführen.
Unter Verwendung der üblichen Visual Studio-Umgebung oder des Standard-gcc-Compilers wird eine Hostanwendung vorbereitet (eine Anwendung vom Typ .exe, die auf dem x64-Hauptprozessor ausgeführt wird), während alle erforderlichen Unterstützungsbibliotheken im SDK enthalten sind. Wenn die Host-Anwendung gestartet wird, wird die FPGA-Firmware geladen, die Daten werden in den Chipkern geladen und die Verarbeitung beginnt gemäß dem konzipierten Algorithmus.

FPGA (FPGA) ist eine vom Benutzer umprogrammierbare, massiv parallele Hardwarestruktur mit Millionen von Logikelementen, Tausenden von DSP-Signalblöcken und Dutzenden von Megabyte Cache für On-Board-Berechnungen, ohne auf die Hauptspeichermodule des Servers zuzugreifen. Mit schnellen E / A-Schnittstellen (10GE, 40GE, 100GE, PCIe Gen 3 usw.) können Sie Daten effektiv mit dem Hauptprozessor des Servers austauschen.


Der OpenCL-Standard ist eine Umgebung zum Ausführen heterogener Software. Die Umgebung besteht aus zwei separaten Teilen:

  1. Host-Software - eine Anwendung, die auf dem Haupt-Zentralprozessor des Servers ausgeführt wird, in C / C ++ geschrieben ist und die OpenCL-API-Funktionen verwendet. Der Host-Server organisiert den gesamten Rechenprozess, liefert die Quell- und Empfangsdaten und interagiert mit allen Serversystemen mit dem FPGA-Beschleuniger.
  2. Accelerator-Software - ein Programm, das in der OpenCL C-Sprache (C-Sprache mit einer Reihe von Einschränkungen) geschrieben und für die Ausführung auf dem FPGA-Chip kompiliert wurde.

Ein typischer Server für paralleles Rechnen ist ein x64-basierter Computer (zum Ausführen von Hostanwendungen), der einen Hardware-FPGA-Beschleuniger enthält, der meistens über den PCI-Express-Bus verbunden ist. Ein solches System wird übrigens im Selectel Lab vorgestellt.

Die Programmier- und Kompilierungssequenz für den FPGA-Beschleuniger besteht aus zwei Schritten. Der Host-Anwendungscode wird von einem Standard-Compiler (Visual C ++, GCC) kompiliert, um eine ausführbare Datei im Server-Betriebssystem (z. B. * .exe) abzurufen. Der Quellcode des FPGA-Beschleunigers (Kernel, Kernel) wird vom AOC-Compiler als Teil des SDK mit dem Empfang einer Binärdatei (* .aocx) vorbereitet. Diese Datei dient nur zur Beschleunigerprogrammierung.

Architektur der OpenCL-Software-Kompilierungsumgebung
Abb. Architektur der OpenCL-Software-Kompilierungsumgebung

Betrachten Sie einen Beispielcode zum Berechnen eines großen Vektors auf zwei Arten
( PS: Schießen Sie nicht auf den Pianisten - im Folgenden wird der Code von der Euler-Projektseite verwendet. )

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,...) ... }

Der Code am Anfang ist ein Beispiel dafür, wie eine Single-Threaded-Implementierung in C mit der Methode der sequentiellen Berechnung von Skalarelementen aussehen könnte.

Die zweite Version des Codes ist eine mögliche Implementierung des Algorithmus auf OpenCL in Form einer Funktion, die auf einem FPGA-Beschleuniger berechnet wird. Es gibt keine Schleife und die Berechnung erfolgt in einer Iteration der Schleife. Die Berechnung eines Vektorarrays erfolgt als Ausführung von N Kopien dieser Funktion. Jede Kopie hat einen eigenen Index, der in einer Schleife in den Iterator eingesetzt wird, und die Anzahl der Wiederholungsversuche wird vom Host festgelegt, wenn der Code ausgeführt wird. Die Iteratoraktion wird von der Funktion get_global_id () bereitgestellt, die mit einem Index innerhalb von 0 ≤ index <N arbeitet.

Kommen Sie zum Punkt: Erstellen eines Fraktals


Die Mandelbrot-Menge ist eine Anordnung von Punkten "c" auf der komplexen Ebene, für die die Wiederholungsrelation Zn + 1 = Zn² + c für Z0 = 0 eine begrenzte Folge definiert.

Wir definieren Zn = Zn + IYn und auch c = p + iq.
Für jeden Punkt wird die folgende Reihenfolge berechnet:

Xn + 1 = Xn² + Yn² + p
Yn + 1 = 2XnYn + q


Die Berechnung der Zugehörigkeit eines Punktes zur Menge bei jeder Iteration wird als Gleichung durchgeführt
Xn² + Yn² <4.

Um das Mandelbrot-Set auf dem Bildschirm anzuzeigen, definieren wir eine Regel:

  1. Wenn die Ungleichung bei einer Iteration gilt, wird der Punkt in die Menge eingegeben und schwarz angezeigt.
  2. Wenn die Ungleichung nicht gilt, beginnend mit einem bestimmten Iterationswert n = N, wird die Farbe durch die Anzahl der Iterationen N bestimmt.

Der Berechnungsprozess auf dem Host ist wie folgt:

  • Die Berechnung der Anzahl der Iterationen für jeden Punkt innerhalb des Pixelfensters wird der Funktion mandel_pixel () zugewiesen.
  • Die sequentielle Aufzählung von Bildpunkten wird von der Funktion softwareCalculateFrame () bereitgestellt. Die Parameter geben das tatsächliche Intervall der berechneten Punkte, den tatsächlichen Schritt des Algorithmus und einen Zeiger auf den Farbpuffer der Bildgröße (theWidth * theHeight) an.
  • Die Farbe des Punktes wird von der SoftColorTable angepasst.

Fahren wir mit dem Code fort:

 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; }

Jedes Pixel wird unabhängig vom anderen berechnet, und daher kann dieser Prozess parallelisiert werden. Bei der Implementierung des Algorithmus für den FPGA-Beschleuniger wird ein SIMD-Befehl erstellt, um die Anzahl für jedes Iterationspixel zu berechnen (Bestimmen des Farbcodes aus der Palette). Die Implementierung von zwei verschachtelten Schleifen im Bildpuffer wird durch Ausführen der Operation (theWidth * theHeight) durch OpenCL gerahmt.

Die Kernelinstanzen in der folgenden Liste werden als Workitem bezeichnet, und die Menge aller Instanzen wird als Indexbereich bezeichnet. Die Merkmale der Hardwarefunktion umfassen Folgendes:

  • Eine Funktionsdeklaration beginnt mit dem Schlüsselwort __kernel.
  • Art der Hardwarefunktion - Die Art des Rückgabewerts ist immer ungültig.
  • Die Rückgabe von Werten erfolgt über Puffer, die als Parameter übergeben werden.
    • Die ersten drei Parameter definieren das Materialgitter, dessen Knoten den Pixeln des Ausgabebildes entsprechen.
    • Der vierte Parameter begrenzt die Anzahl der Iterationen und verhindert so das Schleifen von Punkten, die zum Mandelbrot-Satz gehören.
    • Der fünfte Parameter ist ein Zeiger auf den Ausgabefarbpuffer.
    • Das Schlüsselwort __global gibt den Speichertyp an, über den der Puffer übertragen wird: Dies ist der allgemeine DDR-Speicher (QDR) auf dem Beschleuniger selbst.
    • Das Schlüsselwort "Einschränken" gibt dem Optimierer ein Verbot der Verwendung indirekter Pufferreferenzen.
    • Im 6. Parameter wird ein Zeiger auf die Palette übergeben.
    • Das Schlüsselwort __constant optimiert die Pufferzugriffe, indem ein Cache mit einem schreibgeschützten Attribut generiert wird.

    Die Beschreibung der Funktion in der Liste entspricht in etwa der Implementierung für den x64-Prozessor. Hier erfolgt die Definition der aktuellen Kernelinstanz über die Funktion get_global_id, in die die Dimensionsnummer (0, 1) als Parameter übergeben wird.

    Zur besseren Optimierung wurde eine explizite Angabe des Zyklusbeginns eingeführt. In Ermangelung von Informationen über die Anzahl der Iterationen zum Zeitpunkt der Kompilierung wird die Anzahl der Schleifenschritte explizit angegeben, da für sie eigene Hardwareblöcke erstellt werden. Bei dieser Art der Codierung sollte man aufgrund des FPGA-Ressourcenverbrauchs für eine größere Anzahl von Zyklen auf die Kapazität eines bestimmten auf dem Beschleuniger installierten Chips zurückblicken.

     //////////////////////////////////////////////////////////////////// // 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 #pragma while { int iterations = 0; // Perform up to the maximum number of iterations to solve // the current work-item's position in the image // The loop unrolling factor can be adjusted based on the amount of FPGA // resources available. unroll UNROLL xSqr + ySqr < 4.0 && iterations < maxIterations ) // Perform the current iteration xSqr = x*x; ySqr = y*y; y = 2*x*y + stepPosY; x = xSqr - ySqr + stepPosX; // Increment iteration count iterations++; } // Output black if we never finished, and a color from the look up table otherwise framebuffer[windowWidth * windowPosY + windowPosX] = (iterations == maxIterations) ? BLACK : colorLUT[iterations]; }

    Das Intel FPGA SDK für OpenCL-Dienstprogramm muss vor dem Kompilieren der Hardwareimplementierung des Algorithmus auf dem Host installiert werden. Zu der vorinstallierten Software muss BSP (Board Support Package) des Herstellers der jeweiligen Beschleunigerplatine gehören. In diesem Beispiel wird Intel Quartus Prime Pro 16.1 mit Unterstützung für OpenCL und BSP des Euler Thread Accelerators (Intel Arria 10) installiert.

    Das Folgende ist die Konfiguration von Pfaden und Umgebungsvariablen. Die Variable ALTERAOCLSDKROOT enthält den Pfad zum Intel FPGA SDK, die Variable AOCL_BOARD_PACKAGE_ROOT enthält den Pfad zum Beschleuniger 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\

    Für die Kompilierung wird der aoc-Compiler aus dem SDK verwendet.

     aoc mandelbrot_kernel.cl -o mandelbrot_kernel.aocx --board thread -v -v --report

    Wir entschlüsseln: mandelbrot_kernel.cl - die Datei mit dem Quelltext, mandelbrot_kernel.aocx - die Ausgabeobjektdatei zum Programmieren von FPGA, Thread - den Namen des Beschleunigers aus dem BSP-Paket. Der Schalter --report zeigt einen FPGA-Ressourcennutzungsbericht an. Der Schalter –v zeigt während der Kompilierung Diagnoseinformationen an. Der Ressourcenverbrauchsbericht für den Kernel lautet wie folgt:

    + --------------------------------------------- ------------------- +
    ;; Zusammenfassung der geschätzten Ressourcennutzung;
    + ------------------------------------ + -------- ------------------- +
    ;; Ressource + Nutzung;
    + ------------------------------------ + -------- ------------------- +
    ;; Logiknutzung; 49%
    ;; ALUTs; 26%;
    ;; Spezielle Logikregister; 25%;
    ;; Speicherblöcke; 21%
    ;; DSP-Blöcke; 16%;
    + ------------------------------------ + -------- -------------------;

    Zum Kompilieren der Hostanwendung wurde im Beispiel das Microsoft Visual Studio 2010 Express-Paket mit installiertem Microsoft SDK 7.1 verwendet. In den Projekteinstellungen ist die Konfiguration für x64 ausgewählt. Verbinden Sie anschließend den Ordner für externe Header-Dateien und geben Sie in den Linker-Einstellungen den Pfad zu zusätzlichen Intel FPGA SDK-Bibliotheken an.
    Zusätzliche Verzeichnisse zum Einschließen von Dateien = $ (ALTERAOCLSDKROOT) \ host \ include;
    Zusätzliche Bibliotheksverzeichnisse = $ (AOCL_BOARD_PACKAGE_ROOT) \ windows64 \ lib;
     $(ALTERAOCLSDKROOT)\host\windows64\lib;

    Der allgemeine Aktionsplan zum Starten des Kernels auf dem Beschleuniger lautet wie folgt:

    1. Holen Sie sich eine Liste der Plattformen
    2. Holen Sie sich eine Liste der Geräte
    3. Kontext erstellen;
    4. Laden Sie den Kernel in das Gerät.
    5. Eingabepuffer an das Gerät senden;
    6. Führen Sie den Kernel zur Ausführung aus.
    7. Lesen Sie den Ausgabepuffer vom Gerät.
    8. freier Kontext.

    Betrachten Sie einige Punkte, die in direktem Zusammenhang mit dem Start des Kernels stehen. Ein Kern ist also so konzipiert, dass er ein Pixel des Bildes verarbeitet. Daher müssen Sie N Kernelinstanzen ausführen, wobei N die Gesamtzahl der Pixel im Bild ist.

    Im Folgenden wird der Fall aufgeführt, dass sich mehrere Aufgaben auf dem Server befinden und die Aufgabe auf diese verteilt werden kann. In jedem der Beschleuniger müssen Sie den Kernel laden (Datei mandelbrot_kernel.aocx). Angenommen, die Anzahl der Beschleuniger ist numDevices, und die Bildlinien werden auf alle Beschleuniger aufgeteilt:

     #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); } / / . .

    • Die Funktion createProgramFromBinary erstellt ein OpenCL-Programmobjekt aus einer Objektdatei.
    • Als nächstes wird für jedes Gerät ein Kernel basierend auf dem Programmobjekt erstellt.
    • Die PixelData-Puffer werden erstellt, um die Ausgabe von jedem Kern zu empfangen.
    • Ein Puffer wird zum Speichern der Farbpalette erstellt und in jeden der Beschleuniger geladen.
    • Als nächstes wird für jedes Gerät die Bindung lokaler Anwendungsparameter und Kernelparameter mithilfe der Funktion clSetKernelArg festgelegt.
    • Parameter werden durch Seriennummern in der Kernelfunktionsdeklaration beginnend bei Null bestimmt.

    Der nächste wichtige Punkt ist die Bestimmung der Größe der Aufgabe basierend auf dem Indexbereich gemäß dem globalSize-Array. Dieses Array kann ein-, zwei- oder dreidimensional sein. Für jede Dimension wird eine Dimension als Ganzzahl angegeben. Die Dimension des Space bestimmt die Indexreihenfolge des Workitems im Kernel.

    In dem Beispiel wird für jeden Kern ein zweidimensionaler Raum angegeben, wobei eine der Achsen die Pixelzeilenelemente ist, die zweite der Satz von Bildlinien, die auf diesem Gerät verarbeitet werden. Im Kernelcode wird die Pixelnummer in der Zeile durch Aufrufen von get_global_id (0) erhalten, die Zeilennummer lautet get_global_id (1). Die Variable globalSize wird an die Funktion clEnqueueNDRangeKernel übergeben, um die erforderliche Anzahl der auszuführenden Kernelinstanzen zu starten.

    Nach Abschluss der Ausführung der Kerne werden Pixelpuffer vom Gerät in lokale Arrays gelesen. Lassen Sie uns die Leistung anhand der Anzahl der Bilder pro Sekunde bewerten. Das Ergebnis ist in der Demonstration auf der SelectelTechDay-Konferenz sichtbar ( siehe Anfang des Artikels ).

    Fazit


    Durch die Programmierung von FPGA-Beschleunigern in einer Hochsprache wurde die Zugriffsschwelle für Entwickler auf diese Technologie zweifellos um eine Größenordnung gesenkt. Für diejenigen, die dieses Toolkit gerade beherrschen, gibt es beispielsweise sogar eine FPGA-Implementierung des berühmten Beispiels „Hello World“ .

    Aber nicht so einfach. Das Schreiben - und insbesondere das Debuggen eines klar funktionierenden Algorithmus eines real angewandten Problems erfordert immer noch eine hohe Professionalität. Eine weitere Einschränkung besteht darin, dass jeder FPGA-Chip nur eine Rechenaufgabe innerhalb der Anwendung ausführen kann. Für eine andere Aufgabe muss sie erneut programmiert werden.
    Das Modell der Verwendung der Plattform ermöglicht es Ihnen übrigens, mehr als einen FPGA-Beschleuniger auf dem Host zu haben, obwohl dies eine ziemlich teure Lösung ist.
    Der Host (Hostanwendung) verwaltet den Prozess des Erstellens des Kontexts (Datenstruktur für den Beschleuniger) und der Befehlswarteschlange. Das heißt, Eine einzelne Hostanwendung, in der es verschiedene Unteraufgaben für paralleles Rechnen auf FPGA gibt, kann sie auf verschiedene Beschleuniger laden:
    KERNEL1 => BESCHLEUNIGER A.
    KERNEL2 => BESCHLEUNIGER B.

    Trotzdem lohnt sich der Versuch, FPGA-Beschleuniger zu beherrschen - in vielen Anwendungsbereichen wird diese Technologie unverzichtbar: Telekommunikation, Biotechnologie, Big Data-Verarbeitung, Mustererkennung, Signal- und Bildverarbeitung, Computermathematik und physikalische Feldmodellierung.

    Zusätzliche Informationen zum Artikel:
    www.altera.com ist eine Intel FPGA- Kernressource .
    www.eulerproject.com ist die offizielle Website des Euler-Projekts.
    Altera + OpenCL: Wir programmieren unter FPGA ohne Kenntnis von VHDL / Verilog - ein Artikel über Habr.

Source: https://habr.com/ru/post/de418403/

More articles:


All Articles