Home

inaugural - dissertation - Ruprecht-Karls

image

Contents

1. 6 il we oO Bild 1 1 Vereinfachtes Schaubild von Signalwegen zwischen Genen die Enzyme produzie ren Gen fiinf G5 ist inaktiv der Signalweg zu Enzym B ist unterbrochen somit wird es nicht produziert Zum Zweck einer Untersuchung gibt es die M glichkeit Gene in lebenden Zellen mit einem speziellen Verfahren zu inaktivieren Die Folge ist dass diese Zellen bestimmte Enzyme nicht mehr produzieren Folglich k nnen die Viren die vom speziellen Enzym abh ngig sind sich nicht mehr reproduzieren oder weiterleben und die infizierte Zelle nimmt keinen weiteren Scha den Ben tigen die Viren dieses Enzym nicht wird die Zelle konsumiert und stirbt ab Ein Forschungsgebiet liegt darin den Zusammenhang zwischen Genen und der Enzymproduk tion zu verstehen Meist sind mehrere Gene f r die Produktion einzelner Enzyme verantwort lich Umgekehrt sind aber auch einzelne Gene f r die Produktion mehrerer Enzyme zust ndig Um den Zusammenhang zwischen Genen und Enzymen besser verstehen zu k nnen werden Netzwerke bestehend aus Signalwegen konstruiert die das Wissen veranschaulichen welche Gene f r welche Enzyme zur Produktion ben tigt werden Das Bild 1 1 zeigt eine Unterbre chung des f r die Enzymproduktion wichtigen Zweigs Um die eben genannte Beispielanwendung durchf hren zu k nnen werden mehrere Experten aus unterschiedlichen Gebieten ben tigt Zu ihnen z hlen Biologen Physiker Mathematiker Informatiker und Ingenie
2. get_global_id int D die die Berechnung der globalen Thread Nummer bernimmt wie sie im Quelltextbeispiel 2 1 auf Seite 13 berechnet wurde Quelltextbeispiel 2 2 demonstriert die gleiche Funktion diesmal als OpenCL Kernel __ kernel void CL_matrixAdd __global float A __global float B __global float C int row get_global_id 1 get_global_id 0 int colSize get_global_size 0 int col int element row colSize col C element A element B element Quelltext 2 2 Von vielen Threads durchlaufene OpenCL Kernelfunktion zur Berechnung einer Matrixaddition Wie in CUDA gibt es in OpenCL einen Satz Funktionen Laufzeitumgebung die Kernel Funktionen starten und Daten von und zum Ger te Speicher transferieren Die Laufzeitumge bung von OpenCL umfasst weitere Funktionen zum Anlegen und Ausw hlen eines OpenCL Ger ts sowie die bersetzung einer Kernelfunktion zur Laufzeit Der OpenCL Kompilierer ist kein eigenst ndiges Programm sondern wird durch Bibliotheksfunktionen der Laufzeitum gebung aufgerufen und gestartet D h der OpenCL Programmierer muss erst ein Programm schreiben um seine OpenCL Kernelfunktion bersetzen zu k nnen Der volle Funktionsum fang der Laufzeitumgebung ist in der OpenCL Spezifikation 36 beschrieben Das Bild 2 4 zeigt die OpenCL Strukturen die notwendig sind um ein Programm zu berset zen und zur Ausf hrung zu bringen Die c1 Funktionen generieren die Strukturen bzw ben t
3. Bild 5 9 Vereinfachter SSA AST der Matrixaddition mit zugeordneten Operationen Da jedes Programm Daten aus dem Speicher liest darauf rechnet und Ergebnisse wiederum schreibt befinden sich die Rechenoperationen der Pipeline zwischen den 1oad Operationen und den store Operationen Im Bild ist das lediglich die rote aaa Operation die nacheinander alle Elemente der Matrix addiert Die analysePipelinePart Methode durchsucht den Parse baum beginnend bei den InstTransfer 1oad Operationen vorw rts in Pfeilrichtung zu den InstTransfer store Operationen Mit rekursiven Aufrufen werden alle Verzweigungen durch laufen und alle identifizierten Rechenoperationen als Teil der Pipeline im Bild mit rot mar kiert Die 10ad und die store Operation ben tigt die Angabe an welcher Adresse gelesen bzw geschrieben werden soll Die Adresse wird mit der Instruktion getelemntptr und eventuell weiteren arithmetischen Operationen berechnet die markierte werden sollen Die Markierung hilft die arithmetischen Operationen zur Berechnung der Adresse mit denen aus der Pipeli ne zu unterscheiden Die unterschiedlich markierten arithmetischen Operationen werden sp 78 5 3 VHDL Kompilierer ter in unterschiedlichen VHDL Komponenten berechnet Die Methode analyseTransferPart durchl uft den Parsebaum beginnend bei den InstTransfer 10ad Operationen r ckw rts zu den Kernel Funktionsparametern im Bild zu den blauen Kreisen Beim Durchlauf
4. 28 add nsw 132 26 27 29 load i32 3 align 4 30 getelementptr inbounds i32 29 132 28 store i32 23 i32 30 ret void define void matrixAdd i32 nocapture SA i32 nocapture B i32 nocapture C i32 w nounwind 1 tail call i32 get_ global id i32 0 nounwind 2 tail call i32 get global _id i32 1 nounwind 3 mul nsw 132 2 w 4 add nsw 132 3 1 5 getelementptr inbounds i32 SA 132 4 Mivm opt 6 load i32 5 align 4 7 getelementptr inbounds i32 B i32 34 8 load i32 7 align 4 9 add nsw i32 8 6 10 getelementptr inbounds i32 C 132 4 store i32 9 i32 10 align 4 ret void Bild 5 3 bersetzung der Kernelfunktion mat rixada in die Zwischensprache LLVM IR 69 5 OpenCL zu FPGA bersetzer 5 3 2 Softwarearchitektur VHDL Backend In einem Blockschaubild 5 4 wird die Struktur des VHDL Backend in funktionale Bl cke ge gliedert und anschlie end erl utert Weiter wird mit UML Unified Modeling Language einer grafischen Modellierungssprache die Softwarearchitektur des bersetzers in Klassendiagram men visualisiert VHDL Backend Bild 5 4 Die vier bersetzungschritte des VHDL Backends 1 Ein SSA Parser liest den optimierten LLVM IR SSA Quelltext einer OpenCL Kernel funktion ein und generiert aus den Instruktionen einen Syntaxbaum AST hier SSA AST genannt Wie SSA mit LLVM zusammenh ngt war in 3 3 2 erw hnt Man h t te auch die Speicherr
5. 59 60 61 62 a 63 64 65 66 67 68 69 70 132 Seite 1 6 IEEE 2006 doi 10 1109 FPL 2006 311235 William R Mark R Steven Glanville Kurt Akeley und Mark J Kilgard Cg a system for programming graphics hardware in a C like language ACM Trans Graph Volu me 22 896 907 July 2003 doi 10 1145 882262 882362 Michael McCool Sh Lib URL http libsh org 2003 Microsoft Editor Phoenix Compiler and Shared Source Common Language Infrastruc ture URL https connect microsoft com Phoenix Hubert Nguyen GPU Gems 3 Addison Wesley Upper Saddle River NJ USA 2007 NVIDIA Whitepaper NVIDIA GF100 URL http www nvidia com object 10_86775 html1 2010 NVIDIA Whitepaper NVIDIAs Next Generation CUDA Compute Architecture Fermi URL http www nvidia com content PDF fermi_white_papers NVIDIA_Fermi_Compute_Architecture_Whitepaper pdf 2010 NVIDIA CUDA C Best Practices Guide Mai 2011 NVIDIA NVIDIA CUDA C Programming Guide Mai 2011 NVIDIA CUDA Occupancy Calculator URL http developer download nvidia com compute DevZone docs html C tools CUDA_ Occupancy_Calculator xls 2012 Alexandros Papakonstantinou Karthik Gururaj John A Stratton Deming Chen Jason Cong und Wen Mei W Hwu FCUDA Enabling efficient compilation of CUDA kernels onto FPGAs Application Specific Processors Symposium on Volume 0 35 42 20
6. m ssen In der Doppelschleife werden die Bildmerkmale in der Reihenfolge berechnet welche der Ab h ngigkeitsgraph in Bild 4 3 empfiehlt In den ersten beiden Bl cken werden die Zwischener gebnisse Piy Py und P sowie die Bildmerkmale 4 1 4 5 und 4 9 berechnet Diese Reihenfolge liefert ideale Raten von Cache Treffern in der CPU da stets aus der selben Spei cherregion gelesen wird in der sich die Co Matrix befindet Auch in den folgenden Bl cken bleiben die Speicherzugriffe auf die P und die P _ Speicherregion begrenzt Nur in den letz ten drei Bl cken k nnen die Speicherzugriffe nicht regional begrenzt werden Besonders die Lesezugriffe auf die bereits berechneten Bildmerkmale 4 6 und 4 9 ersparen eine doppelte Berechnung und im Fall von 4 9 wird eine dreifache Berechnung eingespart Die eben vorgestellte Struktur wurde in einer C Klasse implementiert die mit Inline Funk tionen den Funktionsaufruf Overhead einspart Das Programm wurde mit der besten Kompilie reroptimierung bersetzt Das beinhaltet Bem hungen vom Kompilierer die Berechnungen in den Schleifen mit den SSE Instruktionen zu vektorisierten Aus diesem Grund sind die Schlei fen im Quellcode kurz und einfach gehalten 4 3 GPU Implementierung In diesem Kapitel wird dargelegt wie der Haralick Algorithmus auf der GPU Architektur pa rallelisiert wird und es werden Implementierungsdetails dargelegt 4 3 1 Parallele Struktur In CUDA werd
7. Die ersten Grafikkarten waren in ihrer Funktionalit t darauf beschr nkt Text und einzelne farbige Bildpunkte aus dem Grafikspeicher in Videosignale umzusetzen Getrieben von der Computerspieleindustrie wurde in den 90er Jahren die Funktionen um die F higkeit Linien und Fl chen zu zeichnen erweitert um den Grafikaufbau zu beschleunigen Die Grafics Pro cessing Unit GPU war geboren Mit den popul ren 3D Spielen bekam die GPU eine Grafikpipeline die mit mehreren Shader Einheiten aus einer polygonen 3D Rastergrafik ein 2D Bild berechnete Die Shader Vertex Shader Geometry Shader und Pixel Shader konnten programmiert werden um visuelle Ef fekte Objektverformung Lichtquellen mit Schatten Lichtreflexionen und Objekttexturen der 3D Rastergrafik hinzuzuf gen Die zwei g ngigsten Programmierschnittstellen APIs f r die Grafikpipelines sind DirectX 2 von Microsoft und OpenGL 1 von der Khronos Group Die entsprechenden hohen und C hnlichen Shaderprogrammiersprachen sind OpenGL Shading Language GLSL 44 und Microsofts High Level Shading Language HLSL 6 Weiter gibt es die Sprache von NVIDIA C for Grafics Cg 58 die den Einsatz auf Grafikkarten des Herstellers beschr nkt 24 3 1 Co Prozessoren Im Jahr 2003 bertraf die Rechenleistung der damaligen GPUs die der CPUs und setzten sich fortan weiter ab Es war logisch die GPU ebenso f r rechenintensive nicht grafische Algorith men zu nutzen Ein Weg
8. GPU Implementierung 4 3 2 Details der Implementierung Dieser Abschnitt beschreibt Details der Implementierung und Probleme mit deren L sungen Die Beschreibung ist chronologisch zu den Berechnungen beginnt mit dem Speichertranfser in die GPU und endet mit dem Auslesen der Ergebnisse Alle parallelen Operationen werden vereinfacht als singul rer Fall ausgedr ckt Dabei darf nicht vergessen werden dass die Ope rationen parallel auf vielen Matrizen und vielen Zellen aufgef hrt werden 4 3 2 1 Kopie der Zellen Im ersten Schritt wird das Multizellbild in den Arbeitsspeicher der GPU kopiert Nachdem die Zelle mit gegebenen Koordinaten identifiziert wurde werden sie in einen separaten Speicher bereich f r dessen direkten Zugang transferiert Der Zweck des letzten Speichertransfers ist die Erweiterung des Zellbildes mit einem Rahmen der Distanz D 5 Pixel bestehend aus Nullen zum Zellbild Dieser Schritt ist f r die sp tere Pixelpaarbildung notwendig damit auch Pixel am Zellbildrand Paare ber die Grenze des Zellbildes hinaus bilden k nnen Letztendlich wird ein Texturencache auf das Zellbild zum beschleunigten Lesen eingerichtet 4 3 2 2 Lookup Tabellen Es gibt zwei sortierte Nachschlage Tabellen Lookup Tabellen Die erste enth lt Index Grau wertpaare aufsteigend sortiert mit dem Index und die zweite Grauwert Indexpaare aufstei gend sortiert mit dem Grauwert Die Index Grauwerttabelle wird bei der Merkmalsberechnu
9. Werte f r die Mindestmenge des Bustransfers auf Das Bild 5 17 zeigt die Bestandteile des Rechenblocks Sie repr sentieren alle Rechenopera tionen die sich beliebig blockweise verschachteln lassen OPERATION DELAY_INPUT_A BlockComputation DELAY_INPUT_B clkUser rst a_din y Verz a erz gerung Rechen En b_rdy k c_valid Al wer Verz gerung ei b_din Bild 5 17 Schematische Darstellung des VHDL Berechnungsblocks Oben Generic Ports und unten die Entity Ports e Welche Operation das Rechenwerk ausf hren soll wird vom Genericparameter bestimmt Im VHDL Quelltext gibt es f r jede Rechenoperation eine Implementierung die durch den OPERATOR Parameter bestimmt wird 87 5 OpenCL zu FPGA bersetzer e Die DELAY_INPUT Parameter konfigurieren die Anzahl der Registerstufen an jedem Dateneingang f r deren Verz gerung e Das stall Signal unterbricht die Berechnungen D h alle Register halten ihren alten Wert Bild 5 18 zeigt den funktionalen Aufbau des VHDL Blocks der eine bedingte Zuweisung im plementiert BlockCondAssign CMP_FUNCTION CMP_CONSTANT DELAY_INPUT_CMP DELAY_INPUT_B DELAY_INPUT_A clkUser rst cmp_din Vergleicher cmp_rdy a_din Verz gerung b_rdy c_valid Verz gerung b_din stall Bild 5 18 Schematische Darstellung des VHDL Bedingter Zuweisungsblock Oben Generic Ports und unten die Entity Ports e Der Genericparameter CMP_FUNCTION definier
10. der jedoch weiter von der ersten Instanz ben tigt wird In der Funktion systemca11 wird ein v ork verwendet der den Speicherbereich des Vater Prozesses nicht kopiert In den folgenden Abschnitten sind die wichtigsten Passagen der Laufzeitumgebung erl utert wie sie mit dem Rahmendesign der Kommunikationsklasse und letztendlich mit der MPRACE Bibliothek zusammenarbeiten Auf die Fehlerbehandlung wird mit Blick auf das Wesentliche verzichtet Eine vollst ndige und lauff hige OpenCL Anwendung die den Einsatz des Kompi lierers und die Ausf hrung auf dem FPGA demonstriert befindet sich im Anhang A 5 5 2 2 Verwaltung der OpenCL Ger te Der Quelltextausschnitt 5 2 aus der Demonstrationsanwendung im Anhang zeigt wie die OpenCL Funktionen den FPGA findet und f r die weitere Verwendung nutzbar macht cl_int status CL_SUCCESS cl_uint number 0 cl_device_id device NULL cl_context context NULL status clGetDeviceIDs NULL CL_DEVICE_TYPE_ACCELERATOR 1 amp device amp number context clCreateContextFromType NULL CL_DEVICE_TYPE_ACCELERATOR NULL NULL amp status Quelltext 5 2 OpenCL Initialisierung des FPGA als Beschleunigerkarte Die OpenCL Funktion cicetDevice1Ds sucht im Computer ein OpenCL Ger t des Types Be schleuniger ACCELERATOR und gibt eine Liste mit den gefundenen Ger ten zur ck Die Imple mentierung ruft die Call Back Funktion hasDevice auf die ihrerseits eine Instanz von mL
11. einheitlich mit dem gleichen Kompilierer und Optimierer zu verwenden Clang 3 ist ein Open Source Projekt das mehrere Frontends mit dem LLVM Kompilierer framework entwickelt Die Front Ends geh ren zu den Sprachen C C Objektive C und Ob jective C Im Vergleich zu GCC haben Programme die mit einem LLVM Frontend bersetzt wurden eine vergleichbare kompakte Programmgr e besitzen besser optimierten Programm text der sich schneller ausf hren l sst und mit komplexen hohen Sprachkonstrukten ebenso gut 34 3 4 Software Hardware Kompilierer zurecht kommt trotz einfacher Quelltextrepr sentation 3 4 Software Hardware Kompilierer Dieser Abschnitt untersucht mehrere Sprachen welche f r die Biologen gedacht sind um den FPGA Algorithmus in der Offline Prozessierung ver ndern zu k nnen 3 4 1 bersicht Die Designs f r FPGAs haben lange Entwicklungszeiten wegen e der komplexen Beschreibung der algorithmischen Aufgabe in der Register Transfer Ebene e der langen Simulationszeiten e der langen bersetzungszeiten e schwer zu deutender Ubersetzungsfehler e eines m glichen Ungleichverhaltens zwischen Simulationssynthese und Hardwaresyn these e der Schwierigkeit Fehler schnell aufsp ren zu k nnen lange Debug Zyklen Wegen dieser Gr nde existieren sehr viele Software Hardware Kompilierer die die Hardwa resynthese mit C hnlichen Sprachen beschleunigen und erleichtern Im Idealfall muss der
12. male Rechenleistung handelt Die Entwicklungszeiten sind sehr lang Aus diesem Grund gibt es eine F lle von Sprachen die die FPGA Beschreibung vereinfachen k nnen die noch weiter vorgestellt und bewertet werden Vergleichstabelle Zur besseren bersicht werden die Kenndaten in einer Tabelle dargestellt CPU GPU FPGA theoretische Rechenleistung 64 GFLOPS 1500 GFLOPS ca 100 GFLOPS maximale Verlustleistung 90 W 300 400 W 25 W Anschaffungskosten 500 1000 EUR 500 1000 EUR 1000 10000 EUR Entwicklungszeit kurz kurz sehr lang Einsatzfahigkeit sequentiell blockweise gestreamt Tabelle 3 1 Vergleichstabelle der g ngigen Co Prozessoren mit der CPU Die sequentielle Einsatzfahigkeit der CPU wird als Vorteil gewertet denn sie ist in der Lage seriellen Quelltext sehr effizient auszuf hren was die GPU bzw der FPGA nicht gut k nnen Die CPU ist als Universalprozessor ebenfalls in der Lage parallelen Quelltext auszuf hren eine GPU kann dies aber in der Regel effizienter 30 3 2 Haralick Texturen Bildmerkale 3 2 Haralick Texturen Bildmerkale 3 2 1 Beschleunigende Vorarbeiten Es gibt zwei wichtige Arbeiten auf die diese Arbeit aufsetzt Erstens die biologische Anwen dung die die Haralick Bildmerkmale ben tigt 25 und zweitens eine andere die bereits den Haralick Algorithmus in rekonfigurierbarer Hardware beschleunigt hat 71 Die erste Arbeit handelt von der automatischen Analyse
13. ten Operationen w hrend die anderen vom Optimierer entfernt werden Die vollst ndige Optimierung ist erst nach der it Transformation m glich 3 Fertigungsplan Die Analyse findet im Trident IR die Anzahl maximal gleichzeitig ge nutzter Operationen Wenn es f nf Addierer gibt aber nur drei zu beliebigen Zeitpunkten gleichzeitig genutzt werden ben tigt man nun so viele in der Hardware Unter Ber ck sichtigung der verf gbaren FPGA Ressourcen und der verf gbaren Speicherbandbreite wird ein Kontrollflussgraph mit Latenzzeiten erstellt 4 Synthese Einer der ersten Schritte in der Synthese ist die Auswahl der Flie kommabi bliothek die zugrunde gelegt wird Die Ausgabe der Synthese ist ein Hardwaredesign in VHDL Die Synthese umfasst die Generierung der Datenpfade in einer Pipeline dies passiert f r mehrere Bl cke Zu jedem Block wird eine Zustandsmaschine FSM ge heftet die das Zeitverhalten aus dem Kontrollflussgraph ber cksichtigt Die generierten Register dienen der Blockkommunikation untereinander Ein Kontrollmodul steuert die Ausf hrung aller Bl cke Laut der Autoren funktioniert der Trident Kompilierer f r zahlreiche einfache FPGA Anwen dungen Die Leistung des FPGA Designs h ngt ma gebend von der Flie kommabibliothek ab was die Entwickler so wollten Man hat die Wahl Bibliotheken zu verwenden die auf Ressour 38 3 4 Software Hardware Kompilierer cenbedarf oder auf Geschwindigkeit optimiert sind 3
14. 4 1 dargestellt Abb 4 2 a Abb 4 2 b Contrast 4 2 3 625E5 1 035E4 Inverse Different Moment 4 5 0 5558 0 5715 Entropy 4 9 5 5187 5 4807 Tabelle 4 1 Bildmerkmalswerte fiir ein verrauschtes Zellbild 4 2 a und ein weichgezeichne tes b Der Wert fiir das Bildmerkmal Contrast ist gr er f r kontrastreichere Bilder Im eben genann ten Beispiel ist der Wert Contrast f r die verrauschte Zelle tats chlich gr er als der Wert fiir die weichgezeichnete Zelle Umgekehrt ist f r das Bildmerkmal Inverse Different Moment der Wert kleiner f r kontrastreiche Bilder wie in der Tabelle zu erkennen ist Die Entropy ist ein 52 4 1 Untersuchung des Haralick Algorithmusses b Bild 4 2 Zwei gleiche Zellbilder a mit zus tzlichem Rauschen und b weichgezeichnet Ma f r eine Zuf lligkeit deren Wert f r weiche Bilder kleiner ist Weitere Beispiele f r die Anschaulichkeit der Bildmerkmale sind in 73 und 39 zu finden Die Gleichungen der Bildmerkmale aus dem Buch 73 sind f r den allgemeinen Fall symme trische und asymmetrische Co Matrizen Die aufgelisteten Bildmerkmale sind f r symmetrisch quadratische Co Matrizen vereinfacht worden gemeint sind die gleichen Berechnungsergeb nisse f r Zeilen und Spalten Somit konnten in den Bildmerkmalen Correlation 4 3 Infor mation Measure I 4 3 und Definitionen Mean 4 17 Variance 4 18 Entropy 4 18 Terme gek rzt oder durch Ausdr cke m
15. 7 Einschr nkungen in der bersetzung 22 2 22222 90 5 4 Rahmendesign 2 2 2 nun nn 91 3 4 1 Bestandteile 20 2 2 mea don nn a a aa a 91 5 4 2 PCle Core und DMA Engine 2 2 on En 91 543 PCIe Einheit s s s 24 a ae ae RA ed 92 5 4 4 Datenflusskonzept 2 2 2 nn nn 94 5 4 5 Speicherkontroller und vereinfachtes Ansprechen 94 5 4 6 Speicherverwaltungseinheit 2 2 2m on nn 95 5 4 7 Kontrolleinheit 22 on non 98 54 8 Taket e u sd u ae ee a nn a A rca an Be 99 5 5 OpenCL Laufzeitumgebung 2 2 2m a 100 5 5 1 FPGA Kommunikation 2 2 2 22 Co rn nn nn 100 5 5 1 1 Entwicklungssteckkarte ML605 2 222 222 100 5 5 1 2 PCle Treiber und MPRACE Bibliothek 101 5 5 1 3 Kommunikations Klasse 2 22 2 a 102 5 5 1 4 Speichertabellen Klasse 2 2 En nn nn 104 5 5 2 OpenCL Funktionen 2 2 2 m on nn 104 5 5 2 1 Implementierung der OpenCL Funktionen 2 2 2 2 104 5 5 2 2 Verwaltung der OpenCL Ger te 2 2 22 2 nn 105 5 5 2 3 Kernelfunktion bersetzen 2 2 22 2222 106 5 5 2 4 Daten bertragung und Pipeline starten 2 2 2 2 107 5 3 3 Austausch Pipeline Modul 2 22 2m En nn 109 5 5 3 1 Programmierschnittstellen und DPR 2 109 5 5 3 2 Pipeline Module mit DPR austauschen 109 XIII Inhalt 6 Ergebnisse und Diskussion 6 1 GPU Beschleunigung 2 2 2 Cm e 6 2 6 1 1 6 1 2 6 1 3 Gesch
16. Designs Der Ressourcenbedarf wird f r das statische und das dynamischen Design getrennt behandelt Das statische Design umfasst die Logik des Speicherkontrollers der vereinfachten Speicher schnittstelle der PCIe Logik der DMA Engine und dem Rahmendesign Das dynamische De sign entspricht der Pipeline mit der zus tzlichen Logik Da die Pipeline selbst aus Bl cken 115 6 Ergebnisse und Diskussion besteht die unterschiedlich zusammengesetzt werden k nnen werden sie in der Ressourcenta belle 6 4 einzeln aufgelistet Slice Slice Block Slice Registers LUTs RAM FIFO DSP48Els statisches Design 17083 5 7 17245 11 4 34 8 2 Block Addressing 374 0 1 925 0 6 4 0 5 BlockComputing ADD 99 0 0 35 0 0 0 5 BlockComputing MUL 101 0 0 3 0 0 4 BlockCondAssign 132 0 0 43 0 0 BlockTransferRd 2350 0 8 1427 0 9 6 1 4 BlockTransferWr 986 0 3 1562 1 0 4 1 0 Tabelle 6 4 Ressourcenverbrauch des statischen Designs und der Pipelinebausteine Anhand der Ressourcentabelle l sst sich absch tzen wie viele Pipelinebausteine zu einer Pipe line zusammengef gt werden k nnen F r eine hohe Ausnutzung der FPGA Ressourcen muss entsprechend ein gro er Bereich f r das dynamische Design reserviert werden Unter Ausnut zung des gesamten FPGAs einschlie lich des statischen Designs k nnten ca 130 Pipeline bausteine hine
17. Instantiierung sch tzt Beide Sprachen sind hohe Programmiersprachen die es erleichtern Hardware zu synthetisieren Die Synthese beschreibt wie der VHDL Quell text im Hardware Logik Schaltplan bzw in Netzlisten umgesetzt wird Wie das passiert wird in 46 nachvollziehbar gezeigt Ein sehr beliebtes Nachschlagewerk der VHDL Grammatik ist die VHDL Quick Reference Card 26 19 2 Grundlagen Die Sprache VHDL wird f r zwei Zwecke verwendet e Erstens f r die Beschreibung von Logikschaltungen dann spricht man von synthesef higem Quelltext e Zweitens f r die Beschreibung eines Simulationsmodells entsprechend nicht synthese f higem Quelltext Die Unterschiede liegen in Anweisungen die eine zeitliche Abfolge definieren Beispielsweise welche Latenzzeiten eine kombinatorische Logik hat wird von den tats chlich existierenden Logikgattern bestimmt so dass die Angabe einer Zeit nicht synthesef hig ist 2 2 3 Partitionelle Rekonfiguration Die bersetzung der Beschreibung in Hardware passiert in mehreren Schritten Da in die ser Arbeit Xilinx FPGAs verwendet werden wird die bersetzungskette anhand der Xilinx Entwicklungswerkzeuge aus dem System Entwicklungs Nachschlagewerk Development Sys tem Reference Guide 79 vereinfacht demonstriert 1 XST ist das Synthesewerkzeug das wie bereits beschrieben die Hardwarebeschreibung in eine Netzliste bersetzt 2 NGDBuild generiert aus der Netzliste ei
18. Instruktion auf alle Threads ausgef hrt wird m ssen im Fall unter schiedlicher Ausf hrungspfade alle durchlaufen werden D h der Quelltext aller Zweige 2 1 Grafikkarten als Rechenbeschleuniger muss ausgef hrt werden Auch dies gilt es zu vermeiden indem die Daten zusammenge tragen werden die von 32 Threads gleich behandelt werden k nnen Unter dem Begriff divergent branch l sst sich mehr dazu im Benutzerhandbuch finden F r weitere Praktiken der Optimierung sei auf das Benutzerhandbuch verwiesen 2 1 4 Programmiersprache OpenCL OpenCL ist CUDA recht hnlich mit dem Unterschied dass OpenCL keine Begrenzung der eingesetzten Ger te Plattform hat Genau darin liegt der gro e Vorteil in OpenCL geschriebe ne Kernelfunktionen laufen funktionell auf allen GPUs Prozessoren und eingebetteter Hard ware f r die es eine OpenCL Implementierung gibt ohne den Kernel Quelltext ndern zu m ssen Genau wie CUDA verf gt OpenCL ber eine Speicherhierarchie und ein Ausf h rungsmodell welches viele Threads in Gruppen aufteilt die ber gemeinsamen Speicher kom munizieren k nnen und auf der Hardware ausgef hrt werden 15 2 Grundlagen CUDA Begriff OpenCL Begriff deutsche Bedeutung Plattformmodell Hardwarekomponenten device host Stream Multiprozessor Threadprozessor compute device host compute unit CU processing element PE Die Grafikkarte bzw das Ger t Re
19. Screen Wochen f r die Datenaufnahme Mit dem beschleunigten Mikroskop werden nur noch Tage ben tigt was einer ca 20 fachen Beschleunigung entspricht Bild 1 3 Bildausschnitt einer Zellkultur mit hunderten Zellen Es versteht sich von selbst dass eine manuelle Auswertung der Zellbilder bei der enormen Zellenanzahl bzw Datenmenge nicht mehr m glich ist Aus diesem Grund werden Algorith men aus der Bildverarbeitung zur automatischen Analyse verwendet Um m glichst genau mit der manuellen Auswertung bereinzustimmen wurden Algorithmen identifiziert die sich be sonders gut eignen Besonders der Haralick Texturen Merkmal Algorithmus bietet sehr gute Ergebnisse Jedoch besitzt er den Nachteil hoher Rechenintensit t so dass einzelne Compu ter und sogar Rechen Cluster Monate ben tigen um die Bildmerkmale f r alle Zellen eines genomweiten Screens zu berechnen Schon mit dem alten Mikroskop existiert eine gro e Dis krepanz der Aufnahmezeit von Wochen und der ben tigten Zeit der Datenauswertung von Mo naten Die Auswertung der aufgenommenen Daten die auf einem Daten Cluster gespeichert werden wird weiter als Offline Prozessierung bezeichnet W re die Offline Prozessierung ebenso schnell wie die Aufnahmezeit die der Online Prozessierung entspricht k nnte auf die Speicherung der Daten verzichtet werden bzw es werden nur noch die Aufnahmen dauerhaft 1 4 Forschungsfragen gespeichert die von hohem Interesse sind Ve
20. clEnqueueNDRangeKernel queue kernel 1 NULL global_work_size local_work_size 0 NULL NULL status clEnqueueReadBuffer queue mem2 true 0 sizeof unsigned int MEGA h_mem2 0 NULL NULL Quelltext 5 4 Ausschnitt aus der OpenCL Demonstrationsanwendung des Anhangs Kompilierung In der ersten und in der zweiten Zeile wird in der Programm Struktur die einzige Kernelfunk tion extrahiert und eine Kommando Warteschlange erstellt Zeile vier bis acht legt Speicher im Host an und kopiert die Daten in den externen globalen Speicher der ML605 Steckkarte Die Funktion cicreateBuffer greift ber die Call Back Funktion auf die Methode DeviceMem memAllocate zu um einen Speicherbereich zu reser vieren Die Implementierung der Funktion clznqueuewWriteBuffer verwendet die Call Back Funktion devwriteRram um ber einen DMA Kanal die Daten in den externen Speicher zu kopieren Mit Zeile zehn bis zw lf werden die Zeiger der reservierten Speicherbereiche als Parameter der Kernelfunktion bergeben In der c1setkerne1arg Implementierung wird die Zeigeradresse mit der Call Back Funktion devwriteaddress in die Adressblock Komponenten als Basisadres se bertragen Jetzt wei die Pipeline an welcher Adresse die Daten der jeweiligen Kernel Funktionsparameter liegen und sie k nnen vom Adressblock fortlaufend adressiert werden um einen Datenfluss zu generieren In Zeile 14 und 15 wird die Gr e und Dimension der work items
21. der unro11 Anweisung im Clang Kompilierer entfaltet und aufgel st werden Durch die Existenz mehrerer Kernelfunktionen die versetzt auf den Speicher zugreifen kann eine Threadkommunikation erreicht werden Der bisherige bersetzungsprozess ist in der Lage Algorithmen zu bersetzen die Operatio nen auf Werten ohne Interaktion ausf hren Ein Beispiel sind die Intensit tskorrekturen einzel ner Pixel 90 5 4 Rahmendesign FPGA Schnittstellen Rahmendesign statisch Pipelinedesign dynamisch Speicher MEM Speicherverwaltungseinheit kontroller MemManager BA DMA PCle Einheit ee Engine PCleLink m LEDs Do Schalter Bild 5 19 Schnittstellen und Bestandteile des Rahmendesigns 5 4 Rahmendesign 5 4 1 Bestandteile In der Mitte des Bildes 5 19 sind die Bestandteile des Rahmendesigns angeordnet Rechts befindet sich die generierte Pipeline und links sind die FPGA Schnittstellen bzw die FPGA Xilinx Cores mit denen der FPGA verbunden ist Das Rahmendesign bernimmt die Aufgabe Knotenpunkt zwischen Speicher PCIe und Pipeline zu sein und bietet der Software Lauf zeitumgebung eine Kommunikationsschnittstelle ber die sich die Pipeline steuern l sst Die Bestandteile sind folgend weiter erl utert F r das Rahmendesign gibt es keine existierende Implementierung die genutzt werden kann Das Rahmendesign ist speziell f r die FPGA Karte und f r den OpenCL FPGA Kompilierer zu entw
22. dynamische Pipelinemodul platziert wird Au erhalb des Bereichs findet das statische Design die n tigen Ressourcen Bei der Wahl des Bereichs ist es das Ziel dem dynamischen Modul so viel FPGA Chip Fl che wie m glich zu bieten Die Verwendung von DPR bietet an die ICAP Schnittstelle Internal Configuration Access Port zu verwenden Mit ihr kann direkt ber den PCIe Bus ein dynamisches Modul in den FPGA geladen werden ICAP ist mit einer Programmierzeit von Millisekunden schneller als JTAG variabler als der Bootflash und ben tigt kein USB Programmierkabel 5 5 3 2 Pipeline Module mit DPR austauschen Zwischen der dynamischen Pipeline und dem statischen Rahmendesign m ssen Registerstufen existieren damit die Schnittstellensignale einen kurzen Pfad aufweisen ber Modulgrenzen hinweg kann das Zeitverhalten nicht optimiert werden Die betroffenen Signale liegen zwi 109 5 OpenCL zu FPGA bersetzer schen der Pipeline und der Kontrolleinheit im Speicherbus Im Abschnitt 5 3 6 3 wurde bereits darauf hingewiesen welche zus tzliche Logik zur Pipeline ben tigt wird Gemeint sind die Registerstufen im Speicherbus Bevor ein dynamisches Modul im FPGA ausgetauscht wird muss sichergestellt werden dass keine Logik im statischen Design in einen undefinierten Zustand kommen kann W hrend des Austauschs sind die Ausgabesignale des dynamischen Moduls f r kurze Zeit undefiniert und k nnen ein Fehlverhalten im statischen Desig
23. in einer etwas h heren Komplexit t Nach 61 4 Haralick Algorithmus GPU beschleunigt dem die Grauwerte J und J zu den gelesenen Daten der Position i und j nachgeschaut wurden m ssen die Grauwerte subtrahiert und der Absolutwert gebildet werden k Z J Der neu gewonnene Index k ist die Position im Zwischenergebnisvektor P x y k an dem die zuvor gelesenen Daten summiert werden 4 3 2 9 Test und Kontrollimplementierung Neben den Funktionen die den Algorithmus widerspiegeln gibt es weitere f r Testzwecke und f r die Ein Ausgabe Die Aufgaben der zus tzlichen implementierten Funktionen lauten 62 die Laufzeitmessung der einzelnen Kernelfunktionen die Berechnung der Datenraten auf den Grafikkartenspeicher in GBytes s die Berechnung der Rechengeschwindigkeit in GFLOPS das Z hlen des Speicherbedarfs auf der GPU um einen berlauf zu erkennen die Ergebnisausgabe in eine Datei eine Bildschirmausgabe als Fortschrittsanzeige die Fehlersuche um den Inhalt des Grafikkartenspeichers zu kopieren und darzustellen 4 3 GPU Implementierung 4 3 3 Profiling Mit Hilfe des Profilers konnte aus der bisherigen GPU Version eine schnellere GPU Version II implementiert werden Engp sse im Quellcode konnten erkannt sowie weitere Stellen f r Optimierungen identifiziert werden Viele verzweigende Ausf hrungswege und Synchronisa tionspunkte konnten mit nderungen in der Struktur optimiert und
24. k nnen aus diesen Gr nden nicht erreicht werden Unter optimistischen Vor aussetzungen wird in der Praxis eine Rechenleistung von ca 10 GFLOPs 10 von der theo retischen Rechenleistung erreicht Der Unterschied liegt in der geringenen Taktrate und der fehlenden M glichkeit alle Ressourcen f r die Rechnungen einzusetzten 6 2 3 Bandbreite Speicherkontroller Die Speicherbandbreite des DDR3 1066 Speichermoduls hat eine maximale Datenrate von 8 5 GBytes s In diesem Design wird die Benutzerschnittstelle ber die vereinfachte Schnitt stelle mit 133MHz statt 200MHz betrieben In jedem Takt k nnen 32 Bytes bertragen werden Die errechnete Speicherbandbreite liegt bei rund 4 GBytes s Damit die Pipeline nicht unterversorgt wird muss jede Transferleseeinheit mit Daten vom Spei cher versorgt sein Jede Transferleseeinheit 7 4 liefert taktweise 4 Bytes an die Pipeline Die maximale Anzahl von Transferleseeinheiten bestimmt sich zu T 32 Bytes pro Takt vom Speicher rd 6 2 4 Bytes proTrans f erleseeinheit undTakt OF Damit die Pipeline kontinuierlich ein Ergebnis pro Takt zu liefern kann d rfen maximal acht Transferleseeinheiten in der Pipeline verbaut sein Es d rfen auch mehr sein aber dann arbeitet die Pipeline nicht mehr optimal und die Berechnung setzt taktweise aus um auf Daten zu war ten Die gleiche Rechnung gilt ebenso f r die Transferschreibeinheiten um einen Datenstau zu vermeiden 6 2 4 Ressourcenbedarf des
25. niedrigere Taktrate daf r kann die Schaltung beliebig ver ndert werden 26 3 1 Co Prozessoren Der bersichtsbericht aus dem Jahr 2006 76 geht auf den Einsatz von FPGAs in der HPC ein Zu dieser Zeit wurden FPGAs als Co Prozessoren immer h ufiger eingesetzt da ihre Lo gikzellen zahlreicher und komplexer wurden um eine hohe Beschleunigung gegen ber den damaligen CPU zu erzielen FPGAs sind hervorragend f r Logik und Integeroperationen geeignet Aus wenigen Logikzel len lassen sich Addierer Vergleicher und andere arithmetische Integer Operatoren synthetisie ren Viele wissenschaftliche Algorithmen ben tigen Flie kommaberechnung Mit einer arith metischen Analyse die in jedem Teil der Berechnung die notwendige Bitbreite Genauigkeit bestimmt lassen sich die Flie kommaberechnungen in Festkommaberechnungen umwandeln Banerjee 21 konvertiert in seiner Arbeit mit Matlab Flie kommaberechnungen in Festkom maberechnung f r rekonfigurierbare Hardware Dies funktioniert so lange die Zahlendynamik gering bleibt Ein anderer Weg ist es Operatoren f r die Flie kommaberechnungen auf dem FPGA zu im plementieren Beauchamp 22 diskutiert den hohen Ressourcenverbrauch eines Flie komma operators in einem FPGA und verwendet die DSP Slices um einen Multiplikation Addition Operator zu beschreiben Lienhart entwickelte in seiner Dissertation 51 eine Bibliothek mit Flie kommaroperatoren die sich in ihrer Rechengenauigkeit par
26. oo 2 damou 59 4 3 2 Details der Implementierung 2 a a a 59 4 3 2 1 KopiederZellen o o sos s sos sos aworo c ea w ee 59 4 3 2 2 Lookup Tabellen noo aa a 59 4 3 2 3 Gepackte Co Matrix gezielt generieren 60 4 3 2 4 Normalisierte Co Matrix 2 2 a a 60 4 3 2 5 Merkmale erzielen durch Aufsummieren 60 4 3 2 6 Index abh ngige Merkmal Gleichungen 61 4 3 2 7 Zwischenergebnisvektor Px y k 2 22 2200 61 4 3 2 8 Zwischenergebnisvektor Px y k 2 22 2200 61 4 3 2 9 Test und Kontrollimplementierung 62 4 3 3 Profiling s io a a A AL AA 63 Inhalt 5 OpenCL zu FPGA bersetzer 65 Il Konzept u 4 Sn Ee aR a a ee a gw e 65 5 2 bersicht 4 5 5 u 20 e a ARA ehren 66 5 3 VHDL Kompilierer 2 22 2 on on nn 67 5 3 1 bersetzungskette 2 2 Co oo 67 5 3 2 Softwarearchitektur VHDL Backend 70 5 3 3 Parsebaum Generierung 2 2 Cm non nn 73 5 3 4 Parsebaum Analyse 020 000 000004 76 5 3 5 Parsebaum bersetzung 22 22 2 0 ee 79 5 3 5 1 SSA AST zu Block AST 2 2 2 2 nn nen 79 5 3 5 2 Zuordnung Instruktion zum VHDL Block 2 2 2 80 5 3 5 3 Verz gerungen in der Pipeline 81 5 3 6 Parsebaum VHDL Wandlung 2 22 2m on m nn 83 5 3 6 1 Generierung der VHDL Pipeline 2 222 2 20 83 5 3 6 2 VHDL Bl cke als Bausteine 2 2 2 on nenn 84 5 3 6 3 Zus tzliche Logik 2 2 2m on nn 89 5 3
27. vereinfachte Schnittstelle kann mit beliebigen Taktraten angesprochen werden Dualport FIFO verbinden die zwei Takt dom nen und machen den Datentransfer zwischen ihnen m glich Intern werden die Komman dos generiert und abgesetzt ohne dass man es an der vereinfachten Schnittstelle merkt 5 4 6 Speicherverwaltungseinheit Die Speicherverwaltungseinheit Bild 5 21 aus dem Rahmendesign hat die Kontrolle ber den DDR3 Speicher und bietet den anderen Schnittstellen einen Zugriff auf den Speicher an 95 5 OpenCL zu FPGA bersetzer Die ddr3 Verbindung ist ber die Speicherzugriffskomponente an den Speicherkontroller ange schlossen auf die alle Speicherzugriffe geroutet werden Die Signale zur PCle Einheit um dem Host das Lesen und das Beschreiben des Speichers zu erm glichen haben den Pr fix dma Die Transferbl cke aus der Pipeline sind ber einen Lesebus bus_rd bzw bus_ctrl_rd und ber einen Schreibbus bus_wr bzw bus_ctrl_wr an die Speicherverwaltungseinheit angeschlossen 96 e Die Signale ctrl_block_start und end bestimmen den Transfermodus der Speicherver waltungseinheit Die Signalquellen liegen in der Kontrolleinheit bzw in der Pipeline die dar ber Auskunft geben ob die Pipeline gerade am Rechnen ist Es gibt zwei Transfer modi den DMA Transfer und den BUS Transfer Die Hauptaufgabe der Speicherverwaltungseinheit ist den Transferbl cken und der PCIe Einheit Speicherzugriffe zu erm glic
28. wird die Dimension der work group und der work items zur Programmlaufzeit in der Kontrolleinheit des Rahmendesigns gespeichert bzw ver ndert Kontrollregister lesen und schreiben Auf das 32 Bit Kontrollregister im Rahmendesign kann mit den Methoden reaactri und writectri zugegriffen werden Die einzelnen Bits sind im Abschnitt 5 4 7 erl utert Basisadressen setzen Mit der Methode writead ress wird zur Programmlaufzeit die Basis adresse eines Zeigers in ein Register der Adresseinheiten geschrieben Der Zeiger wird der Kernelfunktion bergeben und anhand der Parameterposition einer Adresseinheit in der Pipe line zugeordnet Berechnung steuern Die Methode startcomputation liest und schreibt das Kontrollregister um das Rechenbit zu setzen Die Berechnung startet und das Rechenbit bleibt gesetzt solange die Berechnung l uft was mit der Methode iscomputationDone abgefragt werden kann Der Hostthread pollt auf das Rechenbit bis es sich selbstst ndig r cksetzt dann wenn die Berech nung fertig ist Pipeline Austauschfunktionen Die Methoden setconfigModePR isConfigModePR confi gureIcap und getprid sind f r den Austausch einer Pipeline verantwortlich Auf die Funk tionen wird sp ter im Abschnitt 5 5 3 2 bei der Erl uterung der Laufzeitumgebung genauer 103 5 OpenCL zu FPGA bersetzer eingegangen 5 5 1 4 Speichertabellen Klasse Der globale Speicher wird in der Software verwaltet
29. 0 Gerhard Lienhart Beschleunigung Hydrodynamischer Astrophysikalischer Simulationen mit FPGA Basierten Rekonfigurierbaren Koprozessoren Doktorarbeit Universitat Hei delberg 2004 Mingjie Lin Ilia Lebedev und John Wawrzynek OpenRCL Low Power High Performance Computing with Reconfigurable Devices International Conference on Field Programmable Logic and Applications Volume 0 458 463 2010 doi 10 1109 FPL 2010 93 LLVM LLVM Demo bersetzer Webanwendung URL http 11vm org 2004 Scott A Mahlke David C Lin William Y Chen Richard E Hank und Roger A Bring mann Effective compiler support for predicated execution using the hyperblock SIGMI CRO Newsl Volume 23 45 54 December 1992 doi 10 1145 144965 144998 Guillermo Marcus Acceleration of Astrophysical Simulations with Special Hardware Doktorarbeit University Heidelberg 2011 Guillermo Marcus Wenxue Gao Andreas Kugel und Reinhard Manner The MPRACE framework An open source stack for communication with custom FPGA based accelera tors In 2011 VII Southern Conference on Programmable Logic SPL Volume 605 Seite 155 160 IEEE April 2011 doi 10 1109 SPL 2011 5782641 Guillermo Marcus Gerhard Lienhart Andreas Kugel und Reinhard M nner On Buffer Management Strategies for High Performance Computing with Reconfigurable Hardwa re In 2006 International Conference on Field Programmable Logic and Applications i 131 Literatur 58 a4
30. 09 doi 10 1109 SASP 2009 5226333 A Putnam D Bennett E Dellinger J Mason P Sundararajan und S Eggers CHiMPS A C level compilation flow for hybrid CPU FPGA architectures In Field Programmable Logic and Applications 2008 FPL 2008 International Conference on Seite 173 178 sept 2008 doi 10 1109 FPL 2008 4629927 Daniel J Quinlan ROSE Compiler Support for Object Oriented Frameworks Parallel Processing Letters Volume 10 2 3 215 226 2000 Satnam Singh Computing without Processors Queue Volume 9 50 50 50 63 June 2011 doi 10 1145 1978542 1978558 Literatur 71 M A Tahir A Bouridane F Kurugollu und A A Amira A Amira Accelerating the computation of GLCM and Haralick texture features on reconfigurable hardware In A Bouridane Editor Image Processing 2004 ICIP 04 2004 International Conference on Volume 5 Seite 2857 2860 Vol 5 2004 doi 10 1109 ICIP 2004 1421708 72 David Tarditi Sidd Puri und Jose Oglesby Accelerator using data parallelism to pro gram GPUs for general purpose uses SIGOPS Oper Syst Rev Volume 40 325 335 October 2006 doi 10 1145 1168917 1168898 73 Sergios Theodoridis und Konstantinous Koutroumbas Pattern Recognition Third Edition Academic Press An imprint of Elsevier San Diego CA USA 2006 74 J L Tripp K D Peterson C Ahrens J D Poznanovic und M B Gokhale Trident an FPGA compiler framework for floating point algorithms In Field Progra
31. 2 Microsoft DirectX Developer Center URL http msdn microsoft com de de directx 1995 Just Another Hardware Description Language URL http www jhdl org 1997 Open SystemC Initative Defining and Advancing SystemC Standards URL http www systemc org home 1999 Impuls Accelerated Technologies URL http www impulseaccelerated com 2003 Microsoft Programming Guide for HLSL URL http msdn microsoft com en us library windows desktop bb509635 28v VS 85 29 aspx 2003 AMD s Close to the Metal URL http sourceforge net projects amdctm 2007 clang a C language family frontend for LLVM URL http clang 11vm org 2007 AMD Developer Centrale URL http developer amd com sdks amdappsdk downloads pages default aspx 2009 Apple OS X Lion Technical specifications URL http www apple com macosx specs html 2009 NVIDIA Developer Zone URL http developer nvidia com opencl 2009 IBM developerWorks OpenCL lounge URL https 127 Literatur 13 14 15 16 17 18 19 20 21 22 23 128 www ibm com developerworks mydeveloperworks groups service html communityview communityUuid 80367538 d04a 47cb 9463 428643140bF1 2010 Intel OpenCL SDK Intel Software Network URL http software intel com en us articles vcsource tools opencl sdk 2010 Conformt OpenCL Products URL http www khro
32. 4 2 3 CHIMPS Compiling High level Languages into Massively Pipelined Systems CHiMPS entwickelt von Xilinx und der Universit t Washington 68 vereinfacht HPC Programmierern den Umgang mit FPGAs ohne weitere FPGA Kenntnisse anlernen zu m ssen Die Idee ist es ein FPGA in einen CPU Sockel einer Mehrsockel Hauptplatine zu setzen Als Entlastung der CPU l sst sich C Quelltext in eine Pipeline f r den FPGA bersetzen und ausf hren Der FPGA und die CPU teilen sich den gleichen Adressraum D h der FPGA arbeitet mit gleich hoher Speicherband breite und mit sehr geringer Latenzzeit vergleichbar wie die der CPU Auch Interprozesskom munikation zwischen CPU und FPGA ist sehr effizient m glich Bild 3 4 zeigt ein ANSI C Beispiel das in die Zielsprache CHiMPS Target Language CTL bersetzt wird CTL unterst tzt 42 Assemblerbefehle zu denen jeweils ein VHDL Block exis tiert auf die im letzten Schritt der bersetzung die Befehle abgebildet werden Neben den arithmetischen Instruktionen gibt es wenig andere die den Datenfluss beeinflussen bsp if else Konstrukte bzw or Schleifen Enter foo W X Y Z reg u v s t g int foo int x int x reg fl 2 int y int z dis add x y u 2 sub w Ze v gt Teen 3 add zy Vi s int y w 2 een 4 add y z fl int t u amp y 2 5 add u l E int q t s v 6 sub s v _f2 return q de O Gy EL q Exit foo q a C Quelltext b CTL Instruktionen c VHDL Bl
33. 605 105 aur vn 5 OpenCL zu FPGA bersetzer anlegt Erfolgt beim Anlegen keine Exception existiert die Beschleunigerkarte Beim clcreateContextFromType Aufruf wird eine Kontext Struktur angelegt die den Bezug zum OpenCL Ger t h lt und bei der Verwendung der meisten OpenCL Funktionen mit ange geben werden muss Der Standard sieht vor viele OpenCL Ger te in einen Kontext zusam menzufassen Die vereinfachte Implementierung dieser Arbeit st tzt sich auf einem Kontext mit einem Ger t 5 5 2 3 Kernelfunktion bersetzen In OpenCL wird entweder der Kernelfunktion Quelltext zur Programmlaufzeit bersetzt oder es wird eine bersetzte Kernelfunktions Repr sentation aus einer Datei geladen Der Quelltext abschnitt 5 3 zeigt diesen Vorgang const unsigned char fileContent const size_t fileSize bool binary if binary program clCreateProgramWithBinary context 1 amp device amp fileSize amp fileContent NULL amp status else program clCreateProgramWithSource context 1 const charxx amp fileContent amp fileSize amp status status clBuildProgram program 1 amp device NULL NULL NULL Quelltext 5 3 Ausschnitt aus der OpenCL Demonstrationsanwendung des Anhangs Kompilierung Die Implementierung der Funktion cicreateProgramWithBinary liest eine bereits bersetzte Kernelfunktion ein wenn diese beim Programmstart bergeben wurde Wird dem Programm Quelltext einer Kernefun
34. 7 Fazit und Ausblick Rechenzeit Die Illustration zeigt dass die optimierte Version mit drei Tagen Rechenzeit bereits einen wichtigen Teil der Beschleunigung erreicht Mit dieser Beschleunigung ist es den Biolo gen m glich mehrere Datens tze in einer Woche zu analysieren Der Einsatz der Grafikkarte reduziert die Wartezeit von Tagen auf zwei Stunden bzw 45 Minuten Diese Beschleunigung ist ma gebend daf r dass mehrere Tests am selben Tag durchgef hrt werden k nnen In der jungen Disziplin der Systembiologie mit ihren vielen unerforschten Bereichen ist es au eror dentlich wichtig schnell Ideen und Ans tze auf deren Ergebnisse pr fen zu k nnen um die besten Ans tze aufzusp ren Der Einsatz der GPU tr gt ma gebend zu einem schnelleren Er kenntnissgewinn bei Die Online Prozessierung hat die Anforderung gestreamte Daten in Echtzeit verarbeiten zu k nnen an neue Gegebenheiten anpassbar zu machen und leicht programmierbar zu sein OpenCL kombiniert mit FPGAs auf einer Beschleunigerkarte erf llen alle Anforderungen OpenCL ist eine anerkannte weit verbreitete und leicht zu erlernende Sprache FPGAs werden h ufig f r Echtzeitanwendungen eingesetzt und lassen sich durch Neuprogrammieren mit an deren Anwendungen ersetzen Die in dieser Arbeit demonstrierte OpenCL Entwicklung erf llt den Zweck der Online Prozes sierung aus den Anforderungen Einfache Bildverarbeitungsalgorithmen k nnen in Kernel funktionen umgesetzt werd
35. Algorithmen auf die GPU zu portieren besteht in der Programmie rung der drei Shadereinheiten Die Schwierigkeiten bestanden darin e den Algorithmus an die grafische Pipeline anzupassen Beispielsweise waren die Daten strukturen auf grafische Primitive begrenzt e nur ber die grafischen APIs programmieren zu k nnen Dies erfordert viel Wissen ber die Entwicklung von Grafikanwendungen was gar nicht das Ziel ist e keine Freiheiten zu haben beliebig in den Grafikkartenhauptspeicher schreiben zu d r fen Der Datenfluss einer grafischen Pipeline ist gegeben e die Rechenleistung der Shadereinheiten nicht optimalen nutzen zu k nnen weil die Speicherbandbreite nicht ausreicht Zwei wesentliche Entwicklungen die auf den APIs aufsetzen beseitigten bzw reduzierten die Schwierigkeiten der allgemeinen GPU Programmierung besser bekannt als GPGPU Program mierung General Purpose Grafics Processing Unit Sh und Brook bieten eine einfache Pro grammierumgebung ohne Kenntnisse f r Grafikprogrammierung haben zu m ssen und ohne Einschr nkungen in den Datenstrukturen Ein Backend bersetzt den Quelltext in die Shader sprachen Sh 59 ist eine Metaprogrammiersprache die 2009 von RapidMind abgel st wurde BookGPU 24 stell ein Programmierparadigma vor das aus Streams Kernels und Reduktio nen besteht Mit dem kurzlebigen Projekt Close to Metal 7 Anfang 2007 hatte der Grafikkartenhersteller Ati sp ter AMD neben
36. Furthermore OpenCL is a portable language between several computing architectures If an algorithm writ ten in OpenCL is too complex for the FPGA compiler due to the existing restrictions then a compilation for the GPUs in the offline processing environment is still possible Keywords CUDA Co Processor Compiler DPR FPGA GPGPU haralick texture fea tures hardware syntheses highthoughtput microscopy HPC LLVM OpenCL pipeline generator reconfigurable hardware VHDL IH Kurzbeschreibung Wenn Bilder von einem Mikroskop mit hohem Datendurchsatz aufgenommen werden m ssen sie wegen der gro en Bildmenge in einer automatischen Analyse prozessiert werden Es gibt zwei Ans tze die Offlineprozessierung die Verarbeitung der Bilder auf einem Cluster und die Onlineprozessierung die Verarbeitung des Pixelstroms direkt von den Sensoren F r die Bew ltigung der Bilddaten in der Offlineprozessierung setzt diese Arbeit auf Grafik karten und demonstriert eine Implementierung der Haralick Bildmerkmalerkennung in CUDA Dabei wird der Algorithmus um den Faktor 1000 gegen ber einer CPU L sung beschleunigt Dies erm glicht den Biologen weitere Tests und einen schnelleren Erkenntnisgewinn Die Onlineprozessierung setzt auf FPGAs die sich mit den Sensoren elektrisch verbinden las sen Dabei soll sich der Algorithmus dem Bedarf der Biologen entsprechend ver ndern lassen Diese Arbeit zeigt die Entwicklung eines OpenCL FPGA Kompilierer Protot
37. GA bersetzer aufweisen Das Problem wird in Bild 5 12 verdeutlicht Addr0 TransRdO Verz gerung Latenzsumme yo Datenpfad 9 Addr1 TransRd1 Addr2 TransRd2 Berechnungzeit Addr3 TransWrO Bild 5 12 VHDL Pipeline mit synchronisierten Pfadlaufzeiten Links befinden sich die nummerierten Block TransRd und BlockAddr Komponenten von de nen der Datenstrom gelieft wird An dieser Stelle haben sie in der Summe Null Taktverz ge rung zu sehen an der 0 an den Pfadausg ngen der Komponenten Die BlockComp Kompo nenten O und 1 haben unterschiedliche Berechnungszeiten erkennbar an der Taktzahl in den Klammern BlockComp2 hat das Problem dass an dessen Eing ngen die Daten um einen Takt verz gert ankommen w rden Aus diesem Grund besitzen die Komponenten verz gerbare Da teneing nge deren Daten sich um beliebige Takte verz gern lassen in diesem Fall um einen Takt siehe das blauen K stchen Nicht nur unterschiedliche Berechnungszeiten k nnen un terschiedliche Latenzzeiten auf den Datenpfaden erzeugen sondern auch eine verschachtelte Anordnung der VHDL Bl cke wie es mit dem BlockComp3 der Fall ist Der obere Daten zweig weist bereits eine summierte Latenzzeit von neun Takten auf wegen der erforderlichen Synchronit t muss der untere Datenpfad um neun Takte verz gert werden Der Algorithmus der die Latenzzeiten aller Datenpfade mit Verz gerungen synchronisiert ist mit den Hilfsmetho
38. GB 1 800 Nein Ja lesen schreiben Host Constant 64KB 1 800 Nein Ja nur lesen Host Texture global 1 800 Nein Ja nur lesen Host Tabelle 2 1 Unterschiedliche Speicher der GPU Architektur mit Eigenschaften Daten auf reservierte Bereiche im globalen Speicher ausgelagert werden der sich lokaler Spei cher nennt Shared Speicher Dieser Speicherbereich dient zum Datenaustausch unterschiedlicher Threads innerhalb eines CUDA Blocks Eine Threadkommunikation in unterschiedlichen CUDA Bl k ken ist nicht m glich da sie auf unterschiedlichen SMs ausgef hrt werden k nnen Innerhalb des CUDA Blocks m ssen die Speicherzugriffe auf den Shared Speicher synchronisiert wer den da auch die Ausf hrungsreihenfolge der Threads zuf llig ist Beispiel Jeder Thread wartet am Synchronisationspunkt auf alle anderen damit die Threads die Daten lesen so lange war ten bis die Threads die die Daten schreiben fertig sind Erst mit einer Synchronisation kann die G ltigkeit der Daten garantiert werden Globaler Speicher Dies ist der Massenspeicher ber den die GPU und der Host Daten aus tauschen k nnen Er besitzt die gr te Latenzzeit von 400 800 Taktzyklen und ist mit einem zweistufigen Cache versehen Sind die Daten bereits im Cache vorhanden kann sich die La tenzzeit bis auf einen Taktzyklus reduzieren je nachdem ob die Daten im L1 oder L2 Cache liegen Konstanter Speicher Nur der Host kann den konstanten Speicher schreiben Die Kernelfunk ti
39. Grafikkartenspeicher kopiert werden m ssen Ebenso wird die Grafikkarte erst ab einer gewissen Masse an Daten die blockweise verarbeitet werden effizient D h es muss gewartet werden bis ein Block mit Daten gef llt ist dieser dann bertragen berechnet und wieder zum CPU Speicher zur ck kopiert ist In CUDA gibt es Pipelinetechniken gleichzeitig Daten zu bertragen und Ker nelfunktionen auszuf hren um die Wartezeiten zu reduzieren Je nachdem wie viel Zeit die Echtzeitanforderung er brigt k nnen GPUs eingesetzt werden Die Entwicklungszeiten f r CUDA Programme sind vergleichbar mit denen f r die CPU FPGA Im Gegensatz zu den GPUs verarbeiten FPGAs die Daten als Strom von Einzeleleme neten Vorteile der gestreamten Verarbeitung in einer Pipeline sind dass e die Latenzzeiten kurz sind Sobald Daten bereit liegen k nnen sie verarbeitet werden Dies f hrt zu einer Echtzeitf higkeit da die Ergebnisse in vorhersagbarer Zeit Latenz zeit der Pipeline plus bertragungszeit produziert werden e die Daten f r die Verarbeitung nicht zwischengespeichert werden m ssen im Vergleich zur blockweisen Verarbeitung Die Pipelineverarbeitung ben tigt keinen bzw wenig Zwi schenspeicher wenn man Register nicht zu den Speichern z hlt e die Pipeline parallel arbeitet und im FPGA an neue Rechenanforderungen angepasst wer den kann 29 3 Stand der Technik Wie bereits erw hnt eignen sich FPGAs besonders gut f r I
40. HM i r DH H T al N II N Block RAM DSP Teilst ck 1 0 Block IOB Logikzelle LZ Bild 2 5 Bestandteile des FPGAs e Die Logikzellen sind einzelne Elemente die durch Programmierung unterschiedliche Logikfunktionen beherrschen CLB configureable logic block Meist bestehen sie aus einer LUT look up table was nichts anderes als ein SRAM Speicher ist Der CLB kann als Speicher oder als Logikelement benutzt werden indem ber die Adressierung im Speicher Ausgangswerte einer Wertetabelle abgelegt werden So kann ein Logikgatter imitiert werden das die Addressleitungen als Eing nge und das Datenwort als Ausg nge verwendet 18 2 2 Rekonfigurierbare Hardware e Der Blockspeicher Block RAM ist ein fest auf dem FPGA vorgesehener Speicher um eventuell die kostbarere Ressource der Logikzellen einzusparen Die Block RAMs besit zen zudem zwei Ports die mit unterschiedlichen Taktraten angesprochen werden k nnen e Wegen der Leistung werden DSP Teilst cke DSP slices auf dem FPGA integriert die Aufgaben aus dem Bereich der digitalen Signalprozessoren bernehmen k nnen Sie bestehen aus Hardware Multiplizierer und Addierer die mit einer geringeren Verz ge rung und einer dennoch hohen Taktrate arbeiten Die Verwendung der Addierer oder Multipizierer aus den DSP slices spart viele CLBs ein aus denen die Arithmetik sonst zusammengesetzt werden m sste e Die Ein und Ausgabebl cke IOB sin
41. INAUGURAL DISSERTATION zur Erlangung der Doktorw rde der Naturwissenschaftlich Mathematischen Gesamtfakult t der Ruprecht Karls Universit t Heidelberg vorgelegt von M Sc Markus Gipp aus Mannheim Tag der m ndlichen Pr fung 7 Mai 2012 Online und Offline Prozessierung von biologischen Zellbildern auf FPGAs und GPUs Betreuer Prof Reinhard M nner Prof Holger Fr ning F r meine liebe Sabine Abstract This work is about images from a high throughput microscopy Because of the huge amount of images the analysis has to be processed in an automatic way There are two approaches the offline processing image processing on a computer cluster and the online processing image processing of the streaming data from the sensors To cope with the image data in the offline processing this work uses graphics cards as accele rators and shows an CUDA implementation of the Haralick Texture Features The accelerated version achieves a speed up of around 1000 against a CPU solution This offers the biologist the opportunity to do more tests and leads to a faster gain of knowledge The online processing uses FPGAs which are easy to connect to the sensors The biologists ha ve the constraint to adapt the algorithm for their future needs This work presents a developed OpenCL to FPGA compiler prototype The algorithm can be written in OpenCL and compi led for the FPGA without any knowledge of any hardware description language
42. Mitarbeitern und der entspannten als auch kreativen Atmosph re sehr wohl gef hlt Danke auch an die Korrekturleser Heike Hildenbrand und Thomas Haas sie machten meine S tze dieser Arbeit lesbarer und verst ndlicher Besonderer Dank geht an meine Frau und unsere Familien die mir in der schwierigen Zeit durch die Promotion stets Halt gaben IX Inhalt Abstract Kurzbeschreibung Abk rzungen Inhalt 1 Einf hrung und Ziele des Viroquant Projekts 1 1 Orientierung im Viroquant Projekt 2 2 2m nn nn nn 1 2 Hochdurchsatzmikroskopie 2 22 22 0 nn nn nn 1 3 Beschleunigte Bildverarbeitung 2 2 2m on a 1 4 Forschungsfragen 2 2 2 2 m nn nn nn 2 Grundlagen 2 1 Grafikkarten als Rechenbeschleuniger 2 2 2 Cr m nn nn 2 1 1 Geschichtliche Entwicklung 2 2 2m En m nn 2 1 2 Heutige GPU Architektur 2 2 2 2m on nn 2 1 3 Programmiersprache CUDA 2 En En 2 1 4 Programmiersprache OpenCL o 2 2 Rekonfigurierbare Hardware e 2 2 1 Aufbau eines FPGAs 2 2 on nn 2 2 2 Beschreibungssprache VHDL 2 2 En En nn nn 2 2 3 Partitionelle Rekonfiguration 2 2 2 m on nn 2 3 Kompiliererentwicklung 2 CC nn nn 2 3 1 Frontend soe zu 2a an an A a aboa wog we aos 2 3 2 Backend o d e e we ah a aaa a ws 3 Stand der Technik 3 1 Co Prozessoren s ss sosro nn 3 1 1 Beschleunigung s gt s s s s asaca a asa nn n e e 000008 SANZ GPGPU eo 20a a a a Re
43. Programmierer keine Hardwarekenntnisse haben keine Simulation mehr ausf hren und das Programm bersetzt fehlerfrei mit wenig FPGA Ressourcen und hoher Designfrequenz Mittlerweile gibt es sehr viele C hnliche Sprachen die Entwicklungszeit und die Komplexit t der Hardwarebeschreibungssprachen vermeiden oder verringern Der englische Wikipediaein trag Hardware description language 15 wurde in der Vergangenheit stets aktuell gehalten und beinhaltet eine Liste g ngiger Hardwarebeschreibungssprachen mit Kommentaren ber deren Herkunft bzw deren Eigenschaften Zur Zeit sind 31 Sprachen gelistet darunter die be kannten Sprachen Impulse C von Impuls Accelerated Technologies 5 JHDL von der Brigham Young Universit t 3 und SystemC von ARM CoWare CynApps und Synopsys 4 ber die Liste hinaus gibt es viele weitere Sprachen die global nicht bekannt wurden wie z B CHDL von der Universit t Mannheim 35 3 Stand der Technik Die Sprachen auf die hier weiter eingegangen wird sind in serielle parallele und architek tur bergreifende Sprachen f r den FPGA gegliedert Es ist g ngig Hardware Sprachen auf existierende Software Sprachen aufzusetzen Seriell parallel und architektur bergreifend be zieht sich auf die urspr nglich Sprache 3 4 2 Serielle C Sprachen f r den FPGA 3 4 2 1 Handel C Handel C wurde 2009 von Mentor Graphics in ihr Synthese Softwarepaket aufgenommen Die Sprache ist f r hardwarekundige Progra
44. Re ee TI Vil XIV 23 23 23 24 XI Inhalt 4 XII 3 1 3 FPGA e We au Ale a a dad dd 26 SLA Vergleich a fe ae a 2a Siam a ana hehe 27 3 2 Haralick Texturen Bildmerkale 22 22 on on nen 31 3 2 1 Beschleunigende Vorarbeiten 2 2 a En m rn nen 31 3 2 2 Fazit f r eine Beschleunigung 2 2 KL m rn nn nen 32 3 3 Kompiliererentwicklung 2 Cm nn nn 32 3 3 1 bersicht Kompilierer Bauk sten 22222 222m 33 3 3 2 MALL Mies Son a ie a ioa 33 3 4 Software Hardware Kompilierer 2 2 on on nn 35 3 4 1 bersicht ao e en a 35 3 4 2 Serielle C Sprachen f r den FPGA aoaaa nn nn nen 36 342 1 Handel o e eee Ra Kar de Pe eS 36 S422 TRIDENT coc aa aan wem en 37 34 2 3 CHIMPS 2 2 Con nn 39 3 4 3 Parallele C Sprachen f r den FPGA 40 3431 FEUDA o 2 2a 4 4000 8 e eee nn bebe iss 40 343 2 VOBEnNREL s fee Be Be eo ae en 41 3 4 4 Architekturiibergreifende C Sprachen f r den FPGA 43 344 1 OpenCL cuidas adas Boe Roe Roe ES 43 3 4 4 2 Microsoft Accelerator 2 2 ee 44 3 4 5 Fazit der FPGA Sprachen 2 2 2 on on nen 45 Haralick Algorithmus GPU beschleunigt 49 4 1 Untersuchung des Haralick Algorithmusses 49 4 1 1 Co occurrence Matrizen 2 2 22 m 2 nn 49 4 1 2 Haralick Textur Merkmale 22 Co oo nn 51 4 2 CPU Implementierung 2 2 2 2 a a 55 4 3 GPU Implementierung 2 2 e 55 4 3 1 Parallele Struktur s s Co
45. Soll hei en die Software bestimmt beim Anlegen eines Speicherbereichs welche Adresse vergeben wird Die Singleton Klasse DeviceMem von der es nur eine Instanz gibt speichert alle verwendeten Speicherbereiche mit Adresse und Gr e in einer Liste Mit der Methode bool allocate void pointer unsigned int length wird innerhalb der Liste ein ungenutzter Speicherbereich der entsprechenden L nge gesucht und im Erfolgsfall wird der Zeigerparameter mit einer Adresse gesetzt Die Speicherbereiche werden in der Liste nach der Adresse sortiert um nicht verwendete Speicherbereiche leichter berechnen zu k nnen Der Algorithmus sieht im Fall eines fragmentierten Speichers vor die erste passende L cke zu verwenden Falls f r die angeforderte Speichergr e kein freier Be reich zu finden ist wird ein Nullzeiger und false zur ckgegeben Die Methode bool free void x pointer l scht den zugewiesenen Speicherbereich aus der Liste so dass dieser wieder zur Verf gung steht Weiter besitzt die Klasse DeviceMem eine statische ffentliche Zugriffsmethode zur In stanz eine Methode um die Liste zu l schen um den gesamten Speicher freizugeben und f r Debugzwecke eine Ausgabemethode um alle aktuell belegten Speicherbereiche darzustellen 5 5 2 OpenCL Funktionen 5 5 2 1 Implementierung der OpenCL Funktionen Eine OpenCL Implementierung sieht vor alle Funktionen aus dem Standard 36 zu unter st tzen In dieser Arbeit wurde eine Unter
46. ads existieren die abwechselnd warten und berechnen k nnen Ziel ist es die Auslastung der Be rechnungseinheiten zu optimieren 57 4 Haralick Algorithmus GPU beschleunigt Initialisierungsteil Kernelfunktion OA Kernelfunktion OB Kernelfunktion OC Kernelfunktion OD erstelle Index Grautontabellen setze Co Matrizen auf den Wert Null berechne die Co Matrizen normalisiere die Co Matrizen Teil 1 Lesen von den Co Matrizen Kernelfunktion 1A Kernelfunktion 1B Kernelfunktion IC Kernelfunktion 1D Kernelfunktion 1E Kernelfunktion 1F berechne fi berechne fs berechne fo berechne P berechne P xy berechne Py y Teil 2 Lesen von P Kernelfunktion 2A Kernelfunktion 2B Kernelfunktion 2C berechne mean berechne var berechne H Teil 3 Lesen von P Kernelfunktion 3A Kernelfunktion 3B Kernelfunktion 3C Kernelfunktion 3D x y berechne A berechne fi berechne MacP x y berechne fio Teil 4 Lesen von P Kernelfunktion 4A Kernelfunktion 4B Kernelfunktion 4C berechne fs berechne fg berechne f7 Teil 5 Lesen von den Co Matrizen Kernelfunktion 5A Kernelfunktion 5B Kernelfunktion 5C Kernelfunktion 5D berechne f3 mittels P berechne f4 berechne HXY1 fi2 lese von P berechne HXY2 f 3 lese nur von P Tabelle 4 2 Liste aller Kernelfunktionen in der Ausf hrungsreihenfolge Linke Spalte bein haltet den Funktionsnamen Rechte Spalte beschreibt die Berechnung 58 4 3
47. ag zur besse ren Klassifikation der folgenden linearen Diskriminanzanalyse leistet was eine aufwendige Berechnung unn tig macht Die erstellten Co Matrizen sind mit P gekennzeichnet Alle an deren Definitionen werden weiter unten eingef hrt f h B fa De Js f fo fio i fu f Y k Fsava P y k k 2 2Ng EX Pis y k log Pi y k a Ng hd Pa j 108 Ps j i 1 j Ng 1 Ng 1 L Fr E ey k 0 Ng 1 2 Ps y k log P x y k fo HXY1 H y 1 exp 2 0 HXY2 Fol 4 1 4 2 4 3 4 4 4 5 4 6 4 7 4 8 4 9 4 10 4 11 4 12 4 13 51 4 Haralick Algorithmus GPU beschleunigt Die Definitionen zu den Bildmerkmalen sind in den Gleichungen 4 14 bis 4 21 aufgelistet Ng Ng Prty k LL Punk i j k 2 3 2Ng 2 4 14 i l j Ng Ng P LL el ae aa 4 15 I 1 J Ng Pi L Pij 4 16 J Ng w 8P 4 17 g 1 Ng o P e 8 1 4 18 g 1 Ng Ng HXY1 SE Pos log P 4 19 i l j Ng Ng HXY2 n Po Log PP 4 20 i l j Ng H P g log P e 4 21 g 1 Die meisten der gezeigten Bildmerkmale haben eine visuelle Bedeutung Bild 4 2 zeigt die gleiche Zelle einmal mit zus tzlichem Rauschen a und einmal als weichgezeichnetes Bild b F r beide Zellbilder sind die Bildmerkmale Contrast 4 2 Inverse Different Moment 4 5 und Entropy 4 9 berechnet und in der Tabelle
48. ametrisieren l sst um Hard ware Ressourcen gezielt zu sparen Weitere Entwicklungen von Flie kommabibliotheken fin det man unter 23 2002 und 77 2010 demonstriert an Satellitenaufnahmen mit dem k Means Algorithmus Xilinx bietet heute mit dem Core Generator ein bequemes Programm an Operatoren f r die Flie kommaberechnungen zu erstellen 84 Aufbauend auf seiner Bibliothek entwickelte Lienhart einen Pipelinegenerator 50 Mit einer Sprache bestehend aus mathematischen Ausdr cken lassen sich die Operatoren konfigurieren und zu einer Pipeline verschachteln Bei der Generierung der Pipeline werden automatisch alle n tigen Verz gerungstufen eingesetzt In der Arbeit von Marcus 55 wurde der Pipelinege nerator mit Formeln eines numerischen L sungsverfahrens genannt SPH Smoothed Particle Hydrodynamics demonstriert um astronomische Ereignisse mit gegl tteter Teilchen Hydro dynamik zu simulieren 3 1 4 Vergleich Neben den FPGAs haben sich die GPUs als Co Prozessoren etabliert Welcher der geeignete f r eine Beschleunigung ist h ngt von der Anwendung ab Auch der Cell Prozessor und Pro gramme die mit den SSI Befehlen auf der CPU parallelisiert wurden haben ihren Stellenwert im Beschleunigungssegment Es folgt eine Untersuchung der CPU der GPU und des FPGAs mit Blick auf deren Einsatzf higkeit Energieeffizienz Kosten Rechenleistung und Entwick lungszeit In den folgenden Abschnitten wird auf die Bewertung eingegange
49. asst 384 Bit die kleiner ist als die der Vorg ngerarchitektur Sie besitzt aber wegen der h heren Transferrate von 1848 MHz insgesamt eine h here Datentransferrate von 177 4GBytes s Die Speichertransfers werden von einem Zwei Level Cache System unterst tzt das Zugriffszei ten zum 1 5 GByte gro en Arbeitsspeicher reduziert Der 786 kByte gro e Level Zwei Cache L2 Cache befindet sich neben den Speicherkontrollern auf der Chip Ebene w hrend der L1 Cache sich in den SMs befindet Fr here GPU Architekturen besa en kein Cache System bzw nur eines f r Leseoperationen In der GF100 Architektur ist das Cache System f r Lese und Schreiboperationen ausgelegt und ein Protokoll h lt die drei Stufen L1 L2 und Speicher ko h rent Die Stream Multiprozessoren SM haben eine zentrale Bedeutung in der Architektur Auf ihr werden die Softwareprozesse die aus vielen Threads bestehen k nnen ausgef hrt Die Be standteile eines SMs sind in Bild 2 2 gezeigt 2 1 Grafikkarten als Rechenbeschleuniger 32786 Register 64 kByte geteilter Speicher L1 Cache Bild 2 2 Blockschaltbild eines 16 Stream Multiprozessors aus der Fermi GPU Architektur von NVIDIA 62 Threadprozessoren TP TPs auch als CUDA Cores bezeichnet bestehen aus einer arith metischen Einheit ALU und einer Flie komma Einheit FPU die in mehrere Pipelinestufen zerlegt sind Sie ist die kleinste Recheneinheit die Instruktionen und Operationen ausf hren kan
50. at ihre ei gene Parametrisierung und somit auch ihren eigenen Parseralgorithmus der die Anzahl der Parameter und zus tzlich verwendete Schl sselworte kennt Die meisten Instruktionen spei chern ein Ergebnis in einem neuen Register ab F r jedes Ergebnisregister wird eine Instanz von RegisterValue angelegt die vom value abgeleitet ist Die value Instanzen sind die Zweige 73 5 OpenCL zu FPGA bersetzer im AST die die Instruktionsknoten verbinden linkInstruction findet im AST die Quellregister die von der aktuellen Instruktion verwen det werden und verlinkt sie mit entsprechenden Zeigern Ein Quellregister kann entweder das Ergebnisregister Registervalue einer Instruktion oder ein Parameterregister Parametervalue aus dem Funktionskopf sein dessen Instanzen die sourceattrib Klasse angelegt hatte Wegen der zwei m glichen Quellen wird auf die abstrakte Basisklasse value verwiesen Die sourceParser Klasse bietet viele Suchfunktionen Eintr ge aufsteigend searchInstFwa oder absteigend searchinstRev im SSA AST nach allen m glichen Suchkriterien zu su chen Die Softwarearchitektur ist modular angelegt damit jederzeit weitere Befehle durch Ab leiten der Basisklasse in den Funktionsumfang aufgenommen werden k nnen Spezialisierung LLVM IR Befehl Bedeutung InstCall call Funktionsaufruf zur Bestimmung der Threadindi zes der OpenCL Kernelfunktion InstCmp icmp Vergleich mit parametrisiertem O
51. aten ber den Speicherbus anzufordern die als n chstes in der Pipeline ben tigt werden Die Daten werden in einem FiFo zwischengespei chert aus dem die Pipeline taktweise Datenwerte f r Berechnungen entnimmt BLOCK_NUM BlockTransferRd clkUser rst acalc_fin acalc_addr acalc_rdy acalc_next acalc_burst bus_ctrl_rd_done acalc_valid G ltigkeits bus_rd_burst Puffer bus_ctrl_rd_sel FIFO pus rd addr bus_rd_data bus_rd_addr_valid bus_rd_data_valid data_out bus_rd_data_sel data_rdy data_next data_stall Bild 5 15 Schematische Darstellung des VHDL Transferleseblocks Oben Generic Ports und unten die Entity Ports e Der Genericparameter BLOCK_NUM wird f r den Datenbus zur Block Identifikation ben tigt 85 5 OpenCL zu FPGA bersetzer e Die Zustandsmaschine verfolgt das Ziel den Datenpuffer stets gef llt zu halten ber die acalc Schnittstelle fordert sie eine Adresse an und initiiert ber die bus Schnittstelle einen Speicherzugriff e Nach einer gewissen Latenzzeit kommen die Daten auf mehrere Takte verteilt mit einer Wortbreite von 256 Bit an Es kann oft vorkommen dass nicht alle angeforderten Daten auch von der Pipeline ben tigt werden Deswegen gibt es den G ltigkeitspuffer der alle 32 Bit Werte als g ltige oder als ung ltige Daten maskiert Ung ltige Daten werden nicht in dem Datenpuffer aufgenommen e Die data Schnittstelle erm glicht der Pipeline Daten aus dem Datenpuffer zu entneh me
52. atzhalter und gewinnt an Bedeutung wenn es ein OpenCL Frontend gibt das Shared Speicher unterst tzt 76 5 3 VHDL Kompilierer Sie sollte analysieren wie der Shared Speicher genutzt wird und man k nnte diese Gebrauchs m glichkeiten unterscheiden Der Shared Speicher 1 wird nicht benutzt 2 wird als Puffer f r Speicherzugriffe auf den Hauptspeicher benutzt 3 wird verwendet um Daten darin zu sortieren um lineare Hauptspeicherzugriffe zu er zielen 4 wird f r die Threadkommunikation innerhalb einer work group benutzt Auch ohne den Shared Speicher unterst tzt der jetzige Kompilierer den zweiten Punkt weil in der Pipeline Datenpuffer existieren um eine kontinuierliche Datenversorgung aufrecht zu erhalten analyseBranchPart identifiziert m gliche Verzweigungen im Quelltext e if else Verzweigung Sie ist bisher nicht implementiert Wenn sie implementiert wird m ssen der is Zweig und der eise Zweig gleiche Latenzzeiten gleich viele Instruktio nen besitzen Bei ungleichen Latenzen muss der k rzere Zweig mit zus tzlichen Pipeli nestufen erweitert werden Im Prototyp wird f r die Demonstration lediglich der bedingte Zuweisungs Operator implementiert siehe n chster Punkt e Bedingte Zuweisung Sie ist eine Spezialform einer if e1se Verzweigung und entspricht dem 2 Operator aus C Je nachdem ob ein boolescher Ausdruck wahr oder falsch ist wird entweder der eine oder der andere We
53. beschleunigt werden Speicherzugriffe auf die Matrixzeilen erfolgen gleichzeitig und blockweise entsprechend der Warp Gr e Oft ist die Blockgr e kein Vielfaches der Matrixzeilengr e so dass am Zei lenende Verzweigungen f r die Threads innerhalb eines Blocks entstehen manche Threads lesen und andere warten Die Verzweigungen f hren in der SIMT Architektur zu einer un vermeidlichen seriellen Ausf hrung aller Zweige Damit keine Threads ber das Zeilenende hinaus arbeiten ist es notwendig zu pr fen ob der aktuelle CUDA Block mit seinen Berech nungen bereits an den Zeilenenden der Matrix angekommen ist Diese pr fende i Anweisung konnte f r einen allgemeinen Fall weg optimiert werden Im allgemeinen Fall werden die Ma trixzeilen ohne i Anweisung gelesen ohne den Teil am Zeilenende zu lesen Der danach aus gef hrte Grenzfall beinhaltet die i Anweisung f r die verzweigenden Threads am Zeilenende Diese strukturelle nderung beschleunigt das Leseverhalten der Matrixzeilen weil im h ufige ren allgemeinen Fall keine Verzweigung mehr existiert die ausgewertet werden muss Dieses Konzept wurde in vielen Kernelfunktionen umgesetzt Eine weitere beschleunigende Struktur nderung liegt darin auf den Shared Speicher zu ver zichten Der bliche Weg GPU Daten zu prozessieren ist sie aus dem Grafikkartenspeicher in den Shared Speicher zu lesen anschlie end zu synchronisieren Operationen auf dem Shared Speicher anwenden w
54. boards and kits EK V6 ML605 G htm 2010 133 Literatur 83 Xilinx DDR2 and DDR3 SDRAM Memory Interface Solution 2011 84 Xilinx LogiCORE IP Floating Point Operator v5 0 2011 85 Xilinx Virtex 6 FPGA Memory Interface Solutions 2011 86 Zhiru Zhang Yiping Fan Wei Jiang Guoling Han Changqi Yang und Jason Cong Au toPilot A Platform Based ESL Synthesis System Seite 99 112 2008 doi 10 1007 978 1 4020 8588 8_6 134
55. bt dem Programmierer die Wahl unterschiedliche Flie kommabibliotheken zu verwenden Der bersetzungsprozess ist in vier Phasen unterteilt worden zu sehen in Bild 3 3 1 Das LLVM Frontend bersetzt die Funktionen mit den Berechnungen in LLVM IR Bytecode Der Bytecode ist eine von Menschen nicht lesbare Repr sentation auf die dann ausgew hlte Optimierer angesetzt werden Zuletzt wird der Bytecode in Tridents eigene IR abgebildet bei der nur eine Untermenge an Sprachkonstrukten erlaubt ist Der Trident Kompilierer setzt dabei auf die Arbeit des SeaCucumber Kompilierers 75 auf 37 3 Stand der Technik Trident IR Transformation Zeullngsplaul g Generierung Datenpfade LLVM IR nach Hyperblock Ressourcen Generierung Bytecode Transformation Analyse FSM GCC Front End Auswahl Fertigungsplan Generierung Operationen Auswahl Register Bytecode nach Optimierun Block Generierung Trident IR P 9 Pipelining Kontrollfluss Bild 3 3 bersetzungskette des Trident Kompilierers Optimierung 2 Trident IR Transformation Die Trident IR Repr sentation enth lt zus tzliche Infor mation die die Abbildung der Operatoren Basisbl cke auf Hyperbl cke erlaubt Ein Hyperblock 54 ist die Vereinigung mehrerer Basisbl cke die Kontrollflusssignale nur am Eingang ben tigt S mtliche bedingte it Verzweigungen werden von einer Kontroll abh ngigkeit in eine Dantenabh ngikeit berf hrt Die Auswahl identifiziert die ben tig
56. chen Computer in dem das Rechenger t verbaut ist Skalierbarer Parallel Prozessor mit vielen Rechenkernen Rechenkern als Teil einer gr eren Einheit Speichermodell Speichertypen host memory global memory shared memory local memory host memory global memory local memory private memory Hauptspeicher des Computers in dem das Rechen Ger t ist Hauptspeicher des Rechen Ger ts verbunden mit der CU Geteilten Speicher auf den alle PEs einen CU Zugriff haben Lokal Speicher mit exklusivem Zu griff f r ein PE Ausf hrungsmodell Ausf hrung aller Threads CUDA Block CUDA Thread work group work item Rechen Gruppe die parallel auf einer CU ausgef hrt wird Ausf hrungsfaden als Untermenge ei ner Rechen Gruppe Tabelle 2 2 Fachbegriffe aus CUDA und OpenCL die das selbe meinen Leider existieren unterschiedliche Terminologien in OpenCL zu CUDA die an einer Stelle zu Verwechslung f hren Die Tabelle 2 2 listet die deutsche Bedeutung mit englischen CUDA und OpenCL Fachbegriffen in Bezug auf das Plattform Speicher und Ausf hrungs Modell auf In OpenCL existieren in den Kernelfunktionen keine eingebauten Variablen die den aktu ell arbeitenden Thread identifizieren sondern Funktionen die die work group Nummer bzw die work item Nummer auslesen In OpenCL existiert gegen ber CUDA die Funktion size_t 16 vo Dd WH Bau 2 2 Rekonfigurierbare Hardware
57. chen F r eine effiziente Auslastung der GPU wurde der Haralick Algorithmus auf 26 Kernelfunk tionen herunter gebrochen die einzelne Zwischenergebnisse bzw Bildmerkmale berechnen Die Reihenfolge in der die Zwischenergebnisse und Bildmerkmale berechnet werden l sst sich auch f r die GPU Implementierung vom Abh ngigkeitsgraph Bild 4 3 Seite 54 ablei ten Mit der optimalen Reihenfolge sind die Kernelfunktionen gruppiert worden entsprechend der Leseregion im Grafikkartenspeicher siehe hierzu Tabelle 4 2 Die Gruppierung sorgt f r regional konzentrierte Speicherzugriffe die bei den Texturencaches h here Trefferraten mit resultierendem Geschwindigkeitsgewinn versprechen Wie bereits erw hnt steht f r die Berechnungen jeder individuellen Co Matrix ein Block zur Verf gung Das hei t die Anzahl der Bl cke ist konstant nur die Anzahl der Threads innerhalb der Bl cke muss optimal dimensioniert werden passend zum Ressourcenverbrauch jeder der 26 Kernelfunktionen Die vielen Threads k nnen genutzt werden um gleichzeitig Elemente von der Co Matrix zu lesen und zu berechnen Mit vielen parallelen Threads werden die Speichertransporte durch die Berechnungseinheiten verborgen W hrend ein Teil der Threads auf Daten vom Grafikkartenspeicher warten k nnen andere Threads die Berechnungseinheiten belegen die bereits mit Daten versorgt wurden Das Verstecken der Latenzzeit des Grafikkartenspeichers funktioniert gut sofern gen gend Thre
58. chitektur ab Die Entwicklung ist keine neue Sprache sondern eine Bi bliothek die zu unterschiedlichen Sprachen C C Haskell hinzu gebunden wird Acce lerator erm glicht eine architektur bergreifende bersetzung f r unterschiedliche Zielger te Das Quelltextbeispiel 3 3 zeigt eine eindimensionalen Faltung in der Sprache C die auf meh reren Zielger ten ausgef hrt werden kann using Microsoft ParallelArrays using A Microsoft ParallelArrays ParallelArrays namespace AcceleratorSamples public class Convolver public static float ConvolverlD Target computeTarget float a FloatParallelArray x var n x Length var y new FloatParallelArray 0 0f new n for int i 0 i lt a Length i y a i A Shift x i Hier wird nichts berechnet float result computeTarget ToArraylD y Berechnung auf dem Zielgeraet return result Quelltext 3 3 Beispielquelltext in C der Accelerator Bibliothek 70 In Zeile 10 wird ein FloatParallelArray erstellt das der Speicherung der Ergebnisdaten auf dem Zielger t dient Ebenso wird der Funktion ein paralleles Array bergeben dessen Daten bereits auf dem Zielger t sind Zeile 13 zeigt die Rechenvorschrift der Faltung ohne jegliche Berechnung Der Ausdruck a shift x i wird auf dem Zielger t in eine Speicherzugriffs matrix gewandelt Ist das Zielger t ein FPGA bedeutet dies dass die Daten entsprechend des Schleifenindex an d
59. chreiben und anschlie end den PC neu zu starten um den Link zwischen Bus und PCle Core herzustellen Eine andere M glichkeit besteht in der Verwendung eines Bootflashs das einmal mit einem Design beschrieben wird und beim Bootvorgang das Design in den FPGA l dt Auch dann wird der PCle Core inner halb der Zeitspanne richtig initialisiert Ein weiteres Problem bez glich der Programmierung besteht wenn eine Kernelfunktion in ein Design bersetzt wurde und das FPGA damit programmiert werden soll W hrend des Pro grammierens wird der PCIe Core vom Bus getrennt und die Initialisierung geht verloren Wenn der PC dabei nicht abst rzt ist es zumindest nicht mehr m glich den neu geladenen PCIe Core zu verwenden Die einzige Option den FPGA mit einem neuen Design zu laden ist ein Neu start des PCs Die L sung des Problems besteht darin den initialisierten PCle Core bei einer Programmie rung nicht zu berschreiben und zu halten damit der Link bestehen bleibt Im Allgemeinen ist die Neuprogrammierung lediglich f r den FPGA Bereich notwendig in dem die Pipeline platziert wird Der PCIe Core die Logik des Speicherkontrollers und des Rahmendesigns blei ben f r jede bersetze Pipeline konstant und m ssen somit nicht neu programmiert werden Mit DPR kann dieses Verhalten erzielt werden indem die Pipeline in einem dynamischen Mo dul platziert wird und das restliche Design statisch bersetzt wird Im FPGA wird ein Bereich gew hlt in dem das
60. cke Bild 3 4 bersetzungsverlauf von CHiMPS Der C Quelltext a wird in die Zielsprache b bersetzt und auf VHDL Bl cke c abgebildet Das Bild c zeigt im Datenflussgraph ein FIFO das als Cache von der bersetzung hinzugef gt wurde CHiMPS verfolgt den Ansatz viele kleine lokale verteilte Caches zu verwenden an statt wenige gro e mehrstufige Caches Dies spart FPGA Ressourcen f r die Cache Koh renz ein da jeder lokale Cache einer einzelnen unabh ngigen Funktion zugeordnet ist Um sicher zu stellen dass keine zwei Zeiger auf das selbe Feld zeigen und somit zwei Caches instantiiert werden die inkoh rent werden k nnen muss das ANSI C Schl sselwort restrict verwendet 39 3 Stand der Technik werden Auch wenn es un blicher Programmierstil ist mehrere Zeiger f r das selbe Objekt zu benutzen darf das Schl sselwort nirgendwo fehlen auch wenn man es implizit annehmen k nnte Eine weitere Erg nzung im Quelltext CHiMPS nutzen zu k nnen sind die pragma Anweisungen die den Teil markieren der f r den FPGA bersetzt und ausgef hrt werden soll Mit diesen wenigen nderungen des Quelltextes l sst sich ein Beschleunigungsfaktor von 2 8 bis 36 9 f r bekannte Anwendungen aus der HPC Black Scholes Smith Waterman immul so bol und swm erreichen Es gibt zus tzliche pragma Anweisungen die dem Kompilierer Informationen liefert wie bes sere Rechenleistung zu erwarten ist Ein hardwarekundiger Programmierer mu
61. cke Grafics Processing Unit Grafikprozessor High Performance Computing beschleunigtes Rechnen Intermediate Representation Zwischensprache Low Level Virtual Machine Kompiliererframework Memory Speicher Peripheral Component Interconnect Express Computerschnittstelle Processing Element OpenCL prozessierende Einheit in OpenCL Single Instruction Multiple Data Spezifikation einer Rechnerarchitektur Single Static Assignment einmalige statische Zuweisung Tool Command Language Skriptsprache Very High Speed Integrated Circuit Hardware Description Language Hardwarebeschreibungssprache VII Danksagung Herrn Professor Reinhard M nner danke ich sehr mich an seinem Lehrstuhl aufgenommen und mir eine Promotion erm glicht zu haben Seine Unterst tzung mit seiner menschlichen und freudigen Art machten ihn zu einem idealen Doktorvater Diese Arbeit w re auch nicht wie sie ist wenn mir Guillermo Marcus mit seinem Scharfsinn f r Details keine n tzlichen Tipps f r die Umsetzung gegeben h tte Weiter danke ich ihm und Wenxue Gau deren ausgereifte Entwicklungen PCI Treiber und DMA Logik ich verwenden durfte ohne die meine Arbeit konzeptionell nicht vollst ndig w re Neben den beiden war auch der Austausch mit Andreas Kugel Thomas Gerlach und Nicolai Schn r stets hilfreich wenn es wie h ufig in der Hardwareentwicklung mit den Entwicklungs werkzeugen nicht weiterging Ich habe mich am Lehrstuhl mit all seinen freundlichen
62. d 3 5 zeigt die Aufgabe den Kernel Quelltext mit pragma Anweisungen zu erg nzen Die wichtigsten Anweisungen parametrisieren die parallele Aus f hrung wie das CUDA Grid auf dem FPGA aufgeteilt wird FCUDA GRID und FCUDA 40 3 4 Software Hardware Kompilierer Pragma FCUDA AutoPilot bersetzung CUDA Anweisungen FCUDA AutoPilot Synthes Design Quelltext Quelltext C Quelltext Netzliste Bild 3 5 CUDA FPGA bersetzungskette BLOCKS die Synchronisation welcher Barrieretyp verwendet werden soll FCUDA SYNC die Speicher bertragung wie die Kernelfunktion auf den Speicher zugreifen kann FCUDA TRANSFERS und die Berechnungen wie viele FPGA Rechenkerne verwendet werden FCU DA COMPUTE Im zweiten Schritt der Ubersetzungen wird der FCUDA Quelltext in AutoPilot C konvertiert Die Konvertierung ersetzt die eingebauten parallelen Index Variablen durch Thread Schleifen Dieser Schritt ist vergleichbar mit einer Serialisierung d h die mehrfache Ausf hrung der Kernel Funktion geschieht jetzt mit Hilfe von Schleifen Eine Synchronisation wird mit zwei Teilschleifen ersetzt eine Schleife vor der Synchronisation und eine danach Wenn die erste Teilschleife endet haben alle serialisierten Threads den Synchronisationspunkt erreicht Mit Hilfe der Anweisungen wird die Softwarestruktur hergestellt die AutoPilot f r eine parallele Abarbeitung ben tigt Im dritten Teil der bersetzung wird der AutoPilot Quelltext auf FPGA Softco
63. d Puffer die mit den physikalischen Pins verbun den sind e Das Verdrahtungsnetzwerk basiert auf SRAM Technik was bedeutet dass durch Be schreiben von Speicherzellen die CLBs verdrahtet werden Das hei t auch dass der FPGA beim Spannungsverlust sein Logik Design verliert Wie viel Logik in ein FPGA hineinpasst oder wie viel Chipressourcen ein Logik Design ben tigt h ngt in erster Linie von der Anzahl der CLBs eines FPGAs ab Da aber die F higkeit von CLBs verschiedener Hersteller variieren kann wird f r die Komplexit t der Logikzelle bzw des gesamten FPGAs ein Gatter quivalent Gates verwendet Ein Gatter ist definiert als ein Logikelement mit zwei Eing ngen und einem Ausgang wie beispielsweise bei einem NAND Selbst durch die Angabe aus wie vielen Gates ein CLB FPGA besteht sind die FPGAs unter schiedlicher Hersteller dennoch nicht direkt in Leistung und Ressourcenverbrauch vergleich bar Es gibt viele verschiedene Ma einheiten f r den Ressourcenverbrauch eines Designs bzw f r die Kapazit t eines FPGAs Beispiele sind CLBs LC logic cells slices gates und viele andere 2 2 2 Beschreibungssprache VHDL Um FPGAs zu programmieren wird eine Hardware Beschreibungssprache verwendet Die be kanntesten sind VHDL Very high speed integrated circuit Hardware Description Language und Verilog VHDL ist die erste Wahl in dieser Arbeit weil der Autor die strenge Typpr fung bei der Komponenten
64. d store Befehlen get tigt Weiter verf gt LLVM IR ber einen unendlichen Vorrat an Registern die im Beispiel mit ei nem vorangestellten durchnummeriert werden Jedes Register wird nur einmal beschrieben das der Single Static Assignment Form SSA aus dieser Arbeit 27 entspricht Das Beispiel zeigt weiter dass der Ausdruck b c in IR nur einmal berechnet wird was auf eine erfolgrei che Optimierung schlie en l sst LLVM bietet f nf Vorteile die andere Kompilierer Frameworks nur teilweise haben 1 Das Kompilierermodell h lt die Zwischensprache LLVM IR zu allen Programmlebens zeiten Zeitpunkt des Kompilierens des Linkens zur Ausf hrungszeit und w hrend der Ausf hrungspausen bei Dies erm glicht eine Analyse und Optimierung des Programms in allen Stadien des Programms 2 Neben der Zielsprachengenerierung zur Laufzeit unterst tzt LLVM auch eine rechenin tensive Generierung zur Entwicklungszeit 3 W hrend der Laufzeit kann ein Laufzeitprofil des Programms erstellt werden das zur Optimierung herangezogen wird 4 Das Laufzeitmodell ist transparent Gemeint ist dass die LLVM IR keinem speziellen Objektmodell unterliegt auch keine spezifische Semantik f r Exception verwendet wird und keine bestimmte Laufzeitumgebung ben tigt D h LLVM IR kann f r alle m gli chen Sprachen kompiliert und verwendet werden 5 Wegen der Sprachunabh ngigkeit ist es LLVM m glich alle Quelltexte auch Sprach spezifische
65. das g ngigste und n tzlichste Kompilierer Framework auf das folgend weiter ein gegangen wird 3 3 2 LLVM Low Level Virtual Machine LLVM 48 ist ein Kompilierer Framework das eine st ndi ge Optimierung der zu bersetzenden Programme erm glicht St ndig bedeutet zum Zeit punkt des Kompilierens des Linkens zur Ausf hrungszeit und w hrend der Ausf hrungspau sen Das Kompilierer Framework ist modular aufgebaut so dass ein Austausch des Frontends Quellsprache und des Backends Zielsprache problemlos m glich ist LLVM definiert eine allgemeine Zwischensprache namens Intermediate Representation IR f r die es viele Opti mierer gibt IR ist unabh ngig von einer bestimmten Prozessorarchitektur jedoch einer Assem blersprache recht hnlich int foo int a int b int c return a b c b c Quelltext 3 1 Beispiel C Funktion Als Beispiel ist die C Funktion int f00 in 3 1 gegeben Die bersetzung und Optimie rung in LLVM IR zeigt der Quelltextabschnitt 3 2 Das Beispiel ist mit der Demo Webanwen dung der Seite 53 bersetzt worden 33 3 Stand der Technik define i32 foo i32 a i32 b 132 c nounwind readnone 1 sub nsw 132 b c 2 mul i32 1 a 3 mul i32 2 1 ret i32 3 Quelltext 3 2 Beispiel LLVM IR Die Basisdatentypen in LLVM IR sind Boolean Integer mit w hlbarer Bitbreite Floattypen und Zeiger Die Speicherzugriffe werden ausschlie lich mit den 1oa
66. de sprechen daf r e Der FPGA wird als Kommunikations und Schnittstellenbaustein ben tigt um die Daten vom Sensor in den PC Speicher zu transportieren Da der FPGA sehr leistungsf hig ist kann er auch Co Prozessor Aufgaben bernehmen e Vom Sensor kommen Daten die in den FPGA gestreamt werden Im FPGA werden sie mit geringer Latenzzeit in gestreamter Weise weiterverarbeitet e Die Tatsache dass der Sensor Integerdaten liefert spricht ebenso f r den FPGA Um den Biologen eine Programmierung der Online Prozessierung auf dem FPGA zu erm g lichen wird in dieser Arbeit ein Kompilierer entwickelt der eine Kernelfunktion der Sprache OpenCL in eine Hardwarepipeline in VHDL bersetzt Abschnitt 5 3 Die VHDL Pipeline soll mit umliegender Logik Abschnitt Rahmendesign 5 4 und einer OpenCL Laufzeitumgebung Abschnitt 5 5 auf einer FPGA Karte zur Ausf hrung gebracht werden Folgende Argumente zeigen den Bedarf des Ansatzes e In der Regel beherrschen Biologen kein VHDL sie brauchen eine einfache Sprache die sie f r die Anpassungen der Online Prozessierung einsetzen Handel C ist keine Option 45 3 Stand der Technik da sie lediglich eine vereinfachte Hardwaresprache f r Programmierer mit Hardware kenntnissen ist Auch CHiMPS bedarf komplexer Pragma Anweisungen um die Hard ware zu beschreiben e Im Allgemeinen gilt dass serielle Sprachen entweder zus tzliche Anweisungen vom Programmierer f r eine parallele Aus
67. definiert was der Thread anzahl entspricht Die Gr en werden in die Kontrolleinheit in Register Kopiert Die Adress bl cke ben tigen die Gr en Parallelindizes um das Zugriffsmuster und die letzte Adresse im Datenfluss zu bestimmen Nach der Bekanntgabe der Parallelindizes wird die Berechnung in Zeile 16 gestartet Die Implementierung setzt die Parallelindizes mit der Call Back Funk tion devWriteWorkGroupSize UNd devwriteworkItemsize Die Berechnung wird mit der Call Back Funktion devstartComputation gestartet womit das Bit Null im Kontrollregister gesetzt wird In einer Schleife wird das Bit Null mit devcomputationDone kontinuierlich ausgelesen und gepr ft ob die Berechnung fertig ist Um die CPU Last des aktiven Wartens zu reduzie ren wird innerhalb der Schleife der Hostthread f r wenige Mikrosekunden nach jeder Pr fung schlafen gelegt Die letzte Zeile zeigt wie die Ergebnisse der Berechnung vom externen Speicher des ML605 in den Arbeitsspeicher des Hosts kopiert werden 108 5 5 OpenCL Laufzeitumgebung 5 5 3 Austausch Pipeline Modul 5 5 3 1 Programmierschnittstellen und DPR Der PCIe Bus hat die Anforderung beim Bootvorgang des PCs zu allen PClIe Ger te einen Link herstellen zu m ssen und zwar innerhalb einer Zeitspanne von hundertstel Millisekunden 81 Die JTAG Schnittstelle ist f r diese Anforderung unpraktisch Es besteht die M glichkeit den FPGA mit einem Design inklusive PCIe Core zu bes
68. den 9 10 virtual void writeDMA 101 21 22 23 24 5 OpenCL zu FPGA bersetzer const const const const const const const virtual const const virtual const unsigned int address ES DMABuffers buf unsigned int unsigned int count offset 0 bool inc true Ll bool lock t float timeout void setReg unsigned int unsigned int unsigned int unsigned int Quelltext 5 1 rue 1 0 0 fe address Er value ie getReg ER address FPGA Adresse DMA Puffer Anzahl zu lesender 32Bit Werte Versatz des ersten Wertes Adresse im FPGA inkrementieren Warte bis der Transfer fertig ist Abbruchzeit in Millisekunden FPGA Adresse Wert der geschrieben wird Wert der gelesen wird FPGA Adresse Ausschnitt aus der board h mit den verwendeten Funktionen 5 5 1 3 Kommunikations Klasse Die Kommunikations Klasse hat spezialisierte Methoden die auf den Adressraum des Rah mendesigns siehe Abschnitt 5 4 3 zugeschnitten sind Sie abstrahieren Kommunikations Trans aktionen mit mehreren Aufrufen von Zugriffsfunktionen und verwalten die Klassen Instanzen aus der MPRACE Bibliothek Bild 5 24 zeigt die Kommunikations Klasse mit den abstrahie ren Zugriffsmethoden MPRACE Bibliothek DMABuffer i 0 DeviceTrans writeWorkGroupSize writeWorkltemSize startComputation isComputationDone Bild 5 24 Device Trans Le
69. den Mikroskops ist ein h herer Datendurchsatz von mehreren hundert Megabytes pro Sekunde Die Aufnahmen werden von den Fotosensoren zu einer Ausleseelektronik bertragen Als Ausleseelektronik wird oft eine FPGA Karte in einem PC verwendet Diese hat neben dem Datentransport zum Massenspeicher auch die Aufgabe einer Vorprozessierung der Bilddaten Typische Aufgaben sind Intensit tskorrekturen der Pi xeldaten Sortierung des Pixelstroms bzw die Prozessierung einfacher Bildalgorithmen zur Qualit tsanalyse Dabei existiert die Anforderung das System neuen Gegebenheiten anpassen zu k nnen Wegen des kontinuierlichen Datenstroms muss die Verarbeitung in Echtzeit gesche hen was in dieser Arbeit als Online Prozessierung bezeichnet wird 1 Einf hrung und Ziele des Viroquant Projekts 1 3 Beschleunigte Bildverarbeitung Bei ansteigendem Datenvolumen und ansteigenden Datenraten wird die Speicherung und die Datenauswertung zum Problem Letztendlich ben tigt man immer gr ere Massenspeicher Daten Cluster bei gleich bleibender Geschwindigkeit der Datenauswertung bzw immer leis tungsf higere Rechencluster bei gleich bleibender Speicherkapazit t um die Bilddaten auszu werten F r einen genomweiten Screen wird f r jedes Gen eine Zellkultur Bild 1 3 mit vielen hun dert Zellen aufgenommen Das aufkommende Datenvolumen ist somit ein Produkt hunderter Zellen und tausender Gene Ohne beschleunigtes Mikroskop ben tigt ein genomweiten
70. den aus der sourcecompile Klasse implementiert Der Algorithmus beginnt in der Hilfsmethode 1ayerDetermineRec bei der BlockTransWr Komponente und durchl uft den Block AST r ckw rts mit rekursiven Aufrufen Beim Durchlaufen wird jede Block Komponente mit einer Ebenennummer versehen auf der sie liegt und es wird die gr te Ebene 82 5 3 VHDL Kompilierer bestimmt Im Bild liegt BlockComp3 auf der Ebene null BlockComp2 auf der Ebene eins und die anderen auf der Ebene zwei der gr ten Die Hilfsmethode latencyFillLatencyArray bestimmt f r jede Ebene die gr te Berechnungszeit und speichert sie in einem Array durch die Ebene indiziert Die Arraysumme entspricht dem Datenpfad mit der gr ten Latenz und die Teilsummen Summe der Ebene 0 bis i entsprechen der maximalen Latenz jeder Ebene f r beliebige Datenpfade In der letzten Hilfsmethode 1atencySet InputDelay werden f r jeden Datenteilpfad die tats chlichen Latenzzeiten f r diesen Abschnitt summiert und mit den Teil summen der entsprechenden Ebene verglichen Die Differenz ergibt eine n tige Verz gerung der Eing nge 5 3 6 Parsebaum VHDL Wandlung 5 3 6 1 Generierung der VHDL Pipeline Dieser Abschnitt entspricht dem letzen bersetzungsschritt 4 aus dem Blockschaubild 5 4 Ubersetzungsschritt 3 wird noch erl utert Die vndiarchitecture Klasse aus dem Bild 5 13 er zeugt aus dem Block AST eine VHDL Datei VhdlArchitecture SourceCompiler i i
71. den werden 6 2 2 Nutzen des Pipelinekonzepts Da OpenCL eine parallele Sprache ist ist es einfach eine Kernelfunktion auf den FPGA parallel abzubilden Viele parallel arbeitende Threads durchlaufen die gleichen Operationen eines Pro gramms Es liegt nahe die Operationen eines Threads in eine Pipelinestruktur zu ordnen und sie von vielen Threads durchlaufen zu lassen wie es in dieser Arbeit gemacht wurde Jeder Takt liefert ein Ergebnis f r einen Thread Genau darin besteht der Vorteil des Pipelinekon zepts alle Operatoren in der Pipeline parallel arbeiten zu lassen Die Pipeline ist mit 133MHz getaktet Gibt es in der Pipeline N Operatoren dann ist der Rechendurchsatz R R GigaO perationen s 0 133 GHz N Operationen 6 1 Mit steigendem N und h herer Taktrate wird die Pipeline effizienter Folgende Faktoren lassen den Rechendurchsatz sinken e Die Kernelfunktion besitzt nur wenig Operationen und die Pipeline arbeitet wenig paral 114 6 2 OpenCL Kompilier lel e Die Pipeline ben tigt zu viele FPGA Ressourcen so dass das Zeitverhalten nicht erf llt werden kann und die Taktrate gesenkt werden muss e Bei Pipelines mit vielen Speicherzugriffen k nnte die Speicherbandbreite nicht ausrei chen und die Pipeline muss zeitweise angehalten werden bis alle Teile Transferbl cke wieder mit Daten versorgt sind Die in der Einf hrung vorgestellten theoretisch maximalen FPGA Rechenleistungen von ca 100 GFLOPs
72. der Ethernet bieten bei weitem nicht genug Bandbreite und deren Protokolle werden von den Sensoren nicht unterst tzt Die FPGA Karte soll neben der Datenaufnahme data acquisition DAC ebenso die gestream ten Daten online prozessieren Eine Anpassung des Hardware Designs mit neuen Algorithmen ist in der Regel nicht trivial besonders f r Programmierer ohne Hardwarekenntnisse Aus die sem Grund wurde im Ansatz die Sprache OpenCL ausgew hlt die als einfache Sprache zur Beschreibung der Hardware verwendet werden soll Die Sprache OpenCL bietet in Hinblick auf einen FPGA bersetzer mehrere Vorteile gegen ber anderes Sprachen e OpenCL ist eine parallele Programmiersprache Sie gibt eine Modell vor die Parallelit t in einem FPGA umzusetzen Serielle Programmiersprachen haben diesen Vorteil nicht und es ist schwierig die Parallelit t aus einem seriellen Programm zu extrahieren e Der Hauptvorteil von OpenCL ist die vielseitige Einsatzm glichkeit Mit ihr k nnen unterschiedliche Architekturen programmiert werden wie z B GPUs CPUs und DSPs Die parallele Programmiersprache CUDA ist speziell f r GPUs von NVIDIA entwickelt worden Man k nnte einen CUDA Kompilierer entwickeln der Programme f r FPGAs bersetzt FCUDA ist beispielsweise eine solche Entwicklung aus der Zeit als es noch kein OpenCL gab Warum FCUDA nicht weiter geeignet ist wird im Fazit der FPGA Sprachen erl utert 3 4 5 Es liegt allerdings n her einen Kompilie
73. e nommen Funktionsname GPU Version I GPU Version II Faktor Funktion OA 276 3 ms 275 9 ms 1 Funktion OB 242 4 ms 86 8 ms 2 8 Funktion 0C 466 4 ms 367 6 ms 1 3 Funktion OD 224 4 ms 141 0 ms 1 6 Funktion 1A 221 8 ms 221 2 ms 1 Funktion 1B 416 8 ms 373 6 ms 1 1 Funktion 1C 202 5 ms 203 3 ms 1 Funktion 1D 929 2 ms 177 0 ms 5 2 Funktion 1E 310 1 ms 288 2 ms 1 1 Funktion 1F 602 6 ms 418 2 ms 1 4 Funktion 5A 418 6 ms 300 5 ms 1 4 Funktion 5B 269 3 ms 270 7 ms 1 Funktion 5C 309 9 ms 273 0 ms 1 1 Funktion 5D 225 1 ms 226 5 ms 1 Gesamtausfiihrungszeit inklusive nicht 6600 ms 4650 ms 1 42 gelisteter Funktionen Tabelle 6 3 Ausf hrungszeiten und Beschleunigungsfaktoren der GPU Version I und II mit ausschlie lich relevanten Funktionen 6 2 OpenCL Kompilier 6 2 1 Nutzen der OpenCL Implementierung Mit dieser OpenCL Implementierung ist es wesentlich einfacher den FPGA mit einer Funk tion zu programmieren Beispiel Handelt es sich um eine Vektor Addition Funktion kann 113 6 Ergebnisse und Diskussion diese im FPGA mit einem Addierer und einem Register verwirklicht werden Das ist f r einen hardwarekundigen Programmierer eine leichte Aufgabe und Programmierer ohne Hardware kenntnisse k nnen das lernen Was fehlt ist die Logik die einen Datenfluss aufrecht erh lt und weitere Kontrolllogik die die Berechnungen starten l sst und das Berechnungsende sig nalisiert Die Kommunikation zwischen Hos
74. e zu mehrere Kerne auf einem Chip zu vereinen Eine aktuelle CPU von Intel namens Xeon X7560 mit sechs Kernen und 3 2GHz Taktrate bietet 64GFLOPS das bedeutet 11 GFLOPS pro CPU Kern siehe Intel 18 Die Verlustleistung einer aktuellen CPU liegt bei ca 90W mit einem Preis von ca 500 1000EUR Eine Anwendung die schnell auf der CPU l uft ist beispielsweise sehr IO lastig mit gerin gem Rechenaufwand Solche Anwendungen lassen sich durch Co Prozessoren nicht weiter beschleunigen da die arithmetische Rechenleistung der CPU die Last bew ltigen Eine andere Anwendung w rde aus einem seriellen Algorithmus mit geringem Speicherbedarf bestehen Der kleine Speicherbedarf verspricht eine hohe Trefferrate im CPU Cache Parallele Architek turen bieten keine M glichkeit einer Beschleunigung serieller Algorithmen Die CPU ist ein ausbalanciertes Rechenwerk f r alle Arten von Anwendungen und ist somit das Herzst ck aller Berechnungen Die Co Prozessoren auf Beschleunigerkarten k nnen der CPU einen speziali sierten Teil der Rechenlast abnehmen umgekehrt ben tigt der Co Prozessor eine CPU die ihm eine Aufgabe gibt GPU Die Architektur der GPU wurde bereits im Abschnitt 2 1 2 erl utert Zusammenfassend bietet eine aktuelle GPU 1500GFLOPS hat eine Verlustleistung von 300 400W und einen Preis von ca 500 1000 EUR Das Bild 3 1 zeigt die Leistung verschiedener Prozessoren von NVIDIA und Intel mit deren Rechenleistung und der Speicherbandbr
75. e Berechnungen entlastet In den 90er Jahren als 3D Computerspiele in Mode ka men stieg der Bedarf Rechenoperationen auf die GPU auszulagern Es mussten drei dimen sionale R ume und Figuren mit Polygonen berechnet und im n chsten Schritt mit Texturen gef llt werden Rendern Der stets wachsende Wunsch 3D Spiele immer realistischer werden zu lassen befl gelte die Grafikkartenhersteller immer leistungsf higere GPUs mit immer leis tungsf higere Architekturen zu entwickeln Wann die GPU f r allgemeine Rechenaufgaben als Co Prozessor herangezogen wurde wird sp ter im Stand der Technik Abschnitt GPGPU 3 1 2 behandelt 2 1 2 Heutige GPU Architektur Heutige GPUs haben durch ihre hohe Anzahl von bis zu 512 Prozessorkernen ein vielfaches an Rechenleistung im Vergleich zu CPUs Eine bew hrte Technik hunderte Prozessorkernen auf einem Chip zu vereinen und dessen Komplexit t herunter zu brechen ist die Prozesso renkerne in skalierbar Strukturelemente zu gruppieren die mehrfach auf dem Chip vorhanden sind Dabei sind oft die Strukturelemente wiederum in gleiche Elemente unterteilt Die GPU Architektur die hier beschrieben wird ist die NVIDIA GF100 Architektur 62 die auf dem Vorg nger der Fermi Architektur 63 aufbaut Die Strukturelemente der h chsten Chipebene hei en grafische prozessierende Cluster GPC Vergleichbar mit Multikern CPUs besitzt je der GPC den vollen Funktionsumfang einer GPU Die Elemente de
76. e Slave ID zum Selektieren ermittelt Das Selektionssignal wird synchron mit den Daten und dem G l tigkeitssignal auf den Bus gesendet Die Daten werden von den Datenpuffern des selek tierten Transferblocks aufgenommen 5 4 7 Kontrolleinheit Bild 5 22 zeigt den schematischen Aufbau der Kontrolleinheit aus dem Rahmendesign Die Aufgaben der Einheit sind die Berechnungen der Pipeline zu steuern dem Host Zugriff auf die Register zu geben und die aktuellen Parallelindizes der Pipeline mitzuteilen clkUser rst rstDesign Parallel idx_wr_data Index idx_local_size x y z idx_wr_en Register idx_work_group_size x y z idx_wr_id idx_work_group_id x y z idx_rd_id Kontrollregister idx_rd_data ctrl_wr_reg ctrl_rd_reg ctrl_wr_reg_en id_pipeline ctrl_block_end ctrl_rd_reg_rdy ctrl_pr_config ctrl_block_start Bild 5 22 Schematisches Blockschaltbild der Kontrolleinheit e Die Parallelindizes werden in neun Registern gespeichert die tiber eine Lese und eine Schreibschnittstelle zur PCle Einheit verbunden sind Die Registerinhalte k nnen ber die idx Signale von den Adressbl cken innerhalb der Pipeline ausgelesen werden e Das 32 Bit breite Kontrollregister ist in Kontrollbits Bit 0 bis 15 und in Informationsbits Bit 16 bis 31 aufgeteilt Nur die Kontrollbits k nnen vom Host beschrieben werden Die Kontrollbits steuern die Berechnung den Partieller Rekonfigurations Modus und 98 5 4 Rahmendesign den Pipeline
77. e die Bedingung dass jede Zeilenstartadresse ein vielfaches von 256 sein muss Bei Matrizen die aus mehreren Zeilen bestehen kann am Zeilenende eine L cke mit unbenutzten Elementen entstehen Die Implementierung ber cksichtigt L cken am Zeilenende 4 3 2 4 Normalisierte Co Matrix Die Co Matrizen m ssen f r die Bildmerkmalsberechnung normalisiert sein indem jedes Ele ment durch die Matrixsumme dividiert wird Da die Division eine zeitaufwendige Operation ist multipliziert eine Kernelfunktion alle Matrixelemente mit dem Kehrwert der Matrixsumme Die Ausf hrungszeit wird effizienter weil nur noch eine Division berechnet werden muss 4 3 2 5 Merkmale erzielen durch Aufsummieren Die Gleichungen der Bildmerkmale 4 1 4 13 sowie die Gleichungen der Definitionen 4 14 4 21 bestehen haupts chlich aus Summen Gleichzeitig verf gt jede Kernelfunktion Tabelle 4 2 ber eine ressourcenangepasste Anzahl an parallel arbeitenden Threads die f r die Sum menbildung genutzt werden Die Quelldaten werden blockweise gelesen mit arithmetischen Operationen Multiplikation Logarithmus und andere verrechnet und zu den bereits prozes sierten Daten im Shared Speicher parallel addiert Ergebnis ist ein Block mit Zwischensum men der mit einer parallelen Reduktion zu einer einzelnen Summe addiert wird Die Blockgr e wird definiert durch die Anzahl vorhandener Threads Es gibt zwei Implementierungen der Blockstruktur eine Vektor Block St
78. e die Khronos Group die erste OpenCL Spezifikation 36 eine parallele Sprache die CUDA hnlich ist und die unter anderem f r GPUs und weite re Rechen Architekturen geeignet ist In Abschnitt 3 4 4 1 werden die Implementierungen der Sprache weiter erl utert 3 1 3 FPGA Folgende Liste zeigt Beispiele in wie vielen Bereichen FPGAs eingesetzt werden e Digitale Signalverarbeitung FFTs digitale Signalfilter Kodierer Decodierer CRC Bildverarbeitung Frame Grabber Netzwerktechnik Routen von Datenpakteten mit sehr kurzer Latenzzeit e Glue logic als Kommunikationsbaustein der andere digitale Bausteine miteinander ver bindet e System on a Chip auf dem mehrere funktionelle Einheiten CPU RAM Controller BU Se Peripherie Controller auf einem Chip vereint werden e Data Acquisition DAC Datenaufnahme Vorprozessierung und Weiterverteilung von Sensorwerten e Prototyping im ASIC Entwurf zur kosteng nstigen Entwicklung e High Performance Community HPC Gemeinde zur Beschleunigung von Algorithmen und Rechenanwendungen FPGAs sind n tzliche Prozessoren die beliebige Aufgaben in der digitalen Elektronik ber nehmen k nnen Lediglich bei den HPC Anwendungen wird der FPGA blicherweise in einen PC verbaut in dem er als Co Prozessor fungiert In den anderen Einsatzgebieten sind FPGAs ein preiswerter Ersatz f r ASICs Application Specific Integrated Circuit Der FPGA hat im Vergleich zum ASIC eine deutliche
79. e mit 85 vom Design ver wendet wurden Neuere FPGA Hardware bietet 240 tausend Slices zw lf mal mehr Platz um die fehlenden sechs Bildmerkmale beschreiben zu k nnen und um weitere Beschleunigung zu erziehlen Die zu erwartende Beschleunigung die legt folgenden berlegungen zugrunde die mit Fak toren gesch tzt werden Die damalige Vergleichs CPU Pentium4 mit 2 4 GHz ist circa um den Faktor zwei langsamer als eine heutige CPU Intel 17 mit 3 4 GHz Die FPGA Designfre quenz hat sich seitdem circa um den Faktor zwei erh ht Die fehlenden Bildmerkmale w rden sch tzungsweise 12 Tausend Slices ben tigen da diese aus komplexeren Berechnungen be stehen Das Design w rde dann 30 tausend Slices verwenden die acht mal in einen heutigen FPGA hinein passen w rde Vorausgesetzt es l sst sich ein Konzept entwickeln das diese Par allelit t zul sst Die gesch tzte Beschleunigung eines neuen Designs auf aktueller Hardware 31 3 Stand der Technik verglichen mit einer aktuellen CPU liegt bei circa 50 Der Sch tzwert ermittelt sich aus der damaligen Durchschnittsbeschleunigung der Matrizen und der Bildmerkmale 4 75 7 3 2 und unterliegt folgenden den Faktoren 2 langsamere CPU x2 h here Designfrequenz und x8 mehr FPGA Ressourcen Die Entwicklungszeit das komplexe Design zu berarbeiten und die hohen Kosten f r die FPGA Hardware rechtfertigen die erwartete Beschleunigung nicht Das Ziel einer Beschleu nigung liegt im B
80. eamten Datenelements mit der entsprechenden Threadnummer globale work item berechnet Sie enth lt das Speicherzugriffsmuster das vom sourceAnalyser analysiert wurde und die Basisadresse die von den Kernel Funktionsparametern stammt Eine InstTransfer Klasseninstanz wird immer in eine Kombination von BlockAddr mit BlockTransRd oder BlockTransWr umgesetzt Die Funktion der InstGetzlemptr Klasse r ckt in die Block Addr Komponente Eine bedingte Zuweisung besteht aus einem Vergleich und einer Auswahl der weiter gestreamten Daten vergleichbar mit einem Multiplexer Blocknummer und Typ AddrO TransRdO Addr1 TransRd1 Bild 5 11 Block AST Beispiel der Matrixaddition Addr3 TransWrO Rechenoperation Anhand des Beispiels Matrixaddition soll gezeigt werden wie der Block AST aussieht Bild 5 11 zeigt den entstandenen Pipeline Baum Im Vergleich zum SSA AST 5 7 auf Seite 75 ist dieser einfacher in seinem Aufbau Das liegt daran dass viele Instruktionen auf wenige VHDL Bl cke verteilt wurden Auch die Operationen zur Adressberechnung der Matrixelemente ver bergen sich im Block Addr 5 3 5 3 Verz gerungen in der Pipeline Beim beliebigen Zusammensetzen der VHDL Bl cke zu einer Pipeline k nnen Datenpfade mit unterschiedlicher Latenz entstehen Da am Ende der Pipeline nach der Durchlaufverz gerung jeder Takt ein Ergebnis liefern soll m ssen alle Datenpfade die gleiche Durchlaufverz gerung 81 5 OpenCL zu FP
81. eichertransfers haben und andere die nur einen Bruchteil der maximalen Leistung verwenden In den Spalten Effizienz wurden die erzielte Leistungen im Verh ltnis zu theoretisch maximalen Leistungen dargestellt Es ist zu erkennen dass die Rechengeschwindigkeit von den Speichertransfers gebremst wird da sie eine weitaus h here Ausnutzung hat als die Rechenleistung Weiter l sst sich erkennen dass die Implementierung mit besseren Architekturen skaliert Die Effizienz der Rechenleis tung bleibt im Schnitt bei einem Prozent Der Anstieg auf 1 3 der Recheneffizienz von der GTX280 liegt an der verbesserten Implementierung siehe n chster Abschnitt 112 6 2 OpenCL Kompilier 6 1 3 Optimierungsergebnisse der zweiten Version Die Verbesserungen die im Abschnitt Profiler 4 3 3 auf der Seite 63 beschrieben wurden ergaben ein beschleunigtes Verhalten Die Optimierungskonzepte konnten auf die meisten Ker nelfunktionen angewendet werden In der Tabelle 6 3 sind die Ausf hrungszeiten der zweiten GPU Versionen aufgelistet sowie deren Beschleunigungsfaktor errechnet Im Fokus der Op timierung standen haupts chlich diejenigen Kernelfunktionen mit den l ngsten Ausf hrungs zeiten Die Funktionen 2A 2B 2C 3A 3B 3C und 3D haben jeweils eine Ausf hrungszeit von Ams 6ms Die Funktionen A4 4B und 4C haben Ausf hrungszeiten um die 10ms Diese Funktionen sind bereits sehr effizient wurden nicht optimiert und in die Tabelle nicht aufg
82. einer API auch eine Assemblerschnittstelle f r ihre GPUs bereitge stellt Das er ffnete ebenfalls die M glichkeit Algorithmen auf der GPU zu implementieren CTM wurde 2008 von dem Ati Stream SDK bzw dem AMD FireStream abgel st Mitte 2007 brachte NVIDIA eine Grafikkarte auf den Markt die Vertex Shader Geometry Shader und Pixel Shader in allgemeine Shadereinheiten vereint Das hatte f r die Grafikan wendungen den Vorteil die Rechenlast besser auf die Threadprozessoren verteilen zu k nnen Die Durchlaufrate einer Grafik Pipeline orientiert sich an der langsamsten Einheit w hrend die anderen Einheiten unbesch ftigt bleiben F r die GPGPU Anwendungen hatte die neue Architektur die Flexibilit t hnlich einer CPU Rechenoperationen in beliebiger Reihenfolge ausf hren zu k nnen NVIDIA bot zeitgleich CUDA an womit Programmierer einfach paral lele Programme f r NVIDIAs GPUs entwickelten konnten 25 3 Stand der Technik Mit CUDA konnten bereits viele Algorithmen aus der Forschung beschleunigt werden sie he 61 Typische Beschleunigungen gegen ber einer CPU Version liegen zwischen 3 und 60 Aber auch weitaus h here Werte mehrere Hundert sind in der Literatur zu finden Das Poten zial Algorithmen nur noch in einem Bruchteil der Zeit auszuf hren erm glicht Simulationen mit h herer Aufl sung rechenintensive Algorithmen werden echtzeitf hig bzw erst dann ver wendbar Im Dezember 2008 ver ffentlicht
83. eite zwischen Prozessor und Speicher Viele Algorithmen werden von der verf gbaren Speicherbandbreite gebremst Gerade deswegen ist das Wachstum der Speicherbandbreite ebenso wichtig wie die Rechen 28 3 1 Co Prozessoren GFLOPS GB s A A GeForce GeForce GTX 580 GTX 580 1500 200 GeForce GeForce GTX 480 GTX 480 180 1250 160 GeForce 1000 140 GTX 280 4 NVIDIA GPU GeForce 4 NVIDIA GPU O Intel CPU Se nt 120 O Intel CPU 750 100 GeForce 9800 GTX GeForce 9800 GTX GeForce GeForce 8800 GTX 80 8800 GTX 500 60 Westmere GeForce Westmere 7900 series GeForce 7900 series 250 40 Hapertown Prescott Harpertown Prescott 20 0 0 2003 2004 2005 2006 2007 2008 2009 2010 2011 2004 2005 2006 2007 2008 2009 2010 2011 a Rechenleistung b Speicherbandbreite Bild 3 1 Vergleich CPU und GPU der maximal theoretischen a Rechenleistung und b Band breite 65 leistung Die Grafikkarte eignet sich besonders f r daten parallele Anwendungen mit Daten volumen bis zur Speichergr e der Grafikkarte bei dem jeder Thread unterschiedliche Daten verarbeitet Aber auch task parallele Anwendungen werden von den neueren GPUs unterst tzt indem mehrere Kernelfunktionen zeitgleich gestartet werden F r eine Echtzeitverabeitung eig net sich die GPU nur bedingt da die Daten zwischen CPU Speicher und
84. ellt und in einer Instanzvariablen gespeichert Weitere Attribute die herausgefunden werden betreffen die Quelltextzeilen in welcher der Funkti onskopf die erste und die letzte SSA Instruktion zu finden ist Mit den Zeileninformationen kann die sourceParser Instanz zeilenweise die Instruktionen lesen und einen SSA AST aus Instruction Objekten generieren Abschnitt 5 3 3 Der sourceAnalyser besch ftigt sich mit der Semantik des Programms und analysiert Programmteile wie z B welche Adressierungs art verwendet wird Abschnitt 5 3 4 Mit der Analyse sind alle notwendigen Informationen zusammengetragen um den Block AST aus dem SSA AST zu generieren Abschnitt 5 3 5 Die Instanz vndiarchitecture benutzt den Block AST bestehend aus vna B1ock Instanzen den aufbereiteten Informationen der Analyse und generiert daraus eine VHDL Datei Abschnitt 5 3 6 71 5 OpenCL zu FPGA bersetzer SourceCheck sae lo IN SourceAttrib kx ParameterValue RegisterValue InstructionSet InstructionType lt return gt SourceParser O lt Instruction gt lt param gt SourceAnalyser Parametervalue SourceParser instruction NEUE A AN nstGomp InstComp InstLoad InstStore N Generalisierung abgeleitet von 0 Aggregation SourceCompiler gt lt VhdIBlock gt VhdIPort teil von 4 Komposition teil von Existenzabh ngig VhdlArchitecture VhdlComp Konnektor in Bezie
85. eml if status CL_SUCCESS coutAndExit Could not set kernel argument status status clSetKernelArg kernel 2 sizeof mem2 mem2 if status CL_SUCCESS coutAndExit Could not set kernel argument status const size_t global_work_size 3 El All peo All Shee MEGA 1 1 const size_t local_work_size 3 status clEnqueueNDRangeKernel queue kernel 1 NULL global_work_size local_work_size 0 NULL NULL if status CL_SUCCESS coutAndExit Could not execute kernel status clEnqueueReadBuffer queue mem2 true 0 sizeof unsigned int MEGA h_mem2 NULL NULL if status CL_SUCCESS coutAndExit Could not transfer memory status for i 0 i lt MEGA i Z h_meml i h_mem1 i if Z h_mem2 i 124 0 84 85 86 87 88 89 90 91 92 93 94 95 96 97 break if i MEGA cout lt lt TestProgramKernelExec success lt lt endl else cout lt lt There is a computational miss match at position lt lt i lt lt endl clReleaseMemObject meml clReleaseMemObject mem2 clReleaseCommandQueue queue clReleaseKernel kernel clReleaseProgram program clReleaseContext context Quelltext A 1 Lauff hige OpenCL Anwendung auf dem FPGA 125 Literatur 1 2 LL 3 LL 4 5 6 a 7 8 a 9 10 11 12 Khronos Group OpenGL Standard URL http www khronos org opengl 199
86. en Diese wiederum k nnen f r den FPGA in Pipelinemodule ber setzt werden Dynamisch zur Laufzeit k nnen die Pipelinemodule im FPGA ausgetauscht und ausgef hrt werden Pipeline und Rahmendesign zusammen lassen den Bildverarbeitungsal gorithmus auf dem FPGA ausf hren Das komplexe Zusammenspiel aller Einzelkomponen ten liefert gute Ergebnisse Diese Arbeit erm glicht den Biologen in Bezug auf die Online Prozessierung e die bekannte und weit verbreitete parallele Sprache OpenCL zur Programmierung der FPGA basierten Beschleunigerkarte einzusetzen e ohne Hardwarekenntnisse Kernelfunktionen entwickeln zu k nnen e diese in eine gestreamte Pipeline f r den FPGA zu bersetzen e und die FPGA Beschleunigerkarte mit einfachen Funktionen aus der standardisierten OpenCL Laufzeitumgebung ansteuern zu k nnen Weder FCUDA noch OpenRCL bzw keine bisherige Entwicklung vereint alle aufgelisteten Vorteile Eine zu dieser Arbeit parallele k rzlich erschienene OpenCL Entwicklung f r Altera FPGAs 16 zeigt dass es einen reellen Bedarf gibt FPGAs mit OpenCL programmieren zu k n 120 7 2 Verbesserungen f r die Zukunft nen FPGAs sind sparsamer im Energiebedarf und mit der Pipelinetechnik rechnen sie eine bis zwei Gr enordnungen schneller als Einkern CPUs Der Trend zeigt dass die Taktfrequenz der CPUs kaum noch ansteigt hingegen steigt die Anzahl der Kerne um leistungsf higere CPUs herzustellen Die h here Integrati
87. en Kernelfunktionen auf der GPU vielfach mit unterschiedlichen Threads aus gef hrt Die Threads sind in Bl cken CUDA Bl cken strukturiert die wiederum in einem Grid angeordnet sind siehe in den GPU Grundlagen 2 1 2 D h jeder Thread in jedem Block 55 4 Haralick Algorithmus GPU beschleunigt Schleife ber alle Zellen generiere Matrizen Schleife ber den Winkel A Schleife ber die Distanz D berechne Pxy Pxmy Pxpy berechne fl berechne mean var berechne f2 f10 f11 macPxmy berechne f6 berechne hxyl hxy2 berechne f3 speichere Bildmerkmale l sche Matrizen Bild 4 4 Struktogramm der optimierten Softwareversion des Grids f hrt die gleiche Kernelfunktion auf unterschiedlichen Daten aus Es besteht die Freiheit die angelegten Bl cke ein oder zweidimensional in einem Grid und die Threads drei dimensional innerhalb eines Blocks anzuordnen Dabei entstehen bis zu f nf Indices tx ty tz f r die Threads und bx by f r die Bl cke mit denen jeder Thread nummeriert wird In einer Kernelfunktion werden die Indices genutzt f r jeden Thread unterschiedliche Datenelemente zu adressieren Dabei ist es wichtig die Bl cke und das Grid so zu dimensionieren dass die Indizierung zur Datenstruktur im Speicher passt Die passende Zuordnung der vielen Threads zu den Datenelementen im Speicher erleichtert die Adressierung und bietet lineares Lesen f r schnelle Datentransfers In dieser Anwendun
88. en seiner intensiven Nutzung gut gepflegt d h die Verfah ren werden stets verbessert bzw neue hinzugef gt 68 5 3 VHDL Kompilierer _ kernel void int idx int idy Clidy clang matrixAdd int A int B int C get_global_id 0 get_global_id 1 idx Alidy w idx int w Blidy w idx define void matrixAdd i32 A 132 B i32 SC 132 w nounwind 1 alloca i32 align 4 2 alloca i32 align 4 3 alloca i32 align 4 4 alloca i32 align 4 Sidx alloca i32 align 4 Sidy alloca i32 align 4 store i32 A i32 1 align 4 store i32 B i32 2 align 4 store i32 C i32 3 align 4 store i32 Sw 132 4 align 4 5 call i32 get global _id i32 0 store i32 5 132 idx align 4 6 call i32 get global id i32 1 store i32 6 i32 Sidy align 4 7 load i32 idy align 4 8 load i32 4 align 4 9 mul nsw i32 7 8 10 load i32 idx align 4 11 add nsw 132 9 10 12 load i32 1 align 4 13 getelementptr inbounds i32 12 132 11 14 load i32 13 15 load i32 idy align 4 16 load i32 4 align 4 17 mul nsw i32 15 16 18 load i32 idx align 4 19 add nsw i32 17 18 20 load i32 2 align 4 21 getelementptr inbounds i32 20 132 19 22 load i32 21 23 add nsw i32 14 22 24 load i32 idy align 4 25 load i32 4 align 4 26 mul nsw i32 24 25 27 load i32 idx align 4
89. en sich analog zu der Indizierung eines Matrixelements im linearen Speicher Element Index ix Zeilenlaenge Index j Die globale Threadnummer kombiniert die Threadnummer im Block und die Blocknummer zur Eindeutigkeit Die globale Threadnummer wird f r die Adressie rung des Hauptspeicher benutzt 13 2 Grundlagen CUDA gibt dem Programmierer die Freiheit die Threads und die Speicher ohne Einschr nkun gen zu verwenden zu k nnen Beispielsweise funktioniert es aber nicht effizient wenn einige Threads unterschiedliche Wege im Programm durchlaufen obwohl die Hardware 32 Threads in einem Warp gemeinsam ausf hren muss Es gibt viele Regeln die zu beachten sind um effi ziente Kernel Funktionen zu entwickeln die im Benutzerhandbuch der vorbildlichen Praktiken 64 beschrieben sind Hier ein Auszug der wichtigsten Praktiken 14 Lineare Speicherzugriffsmuster sind am effizientesten da der Speicherkontroller keinen schnelleren Betrieb hat als den Burst Modus Soll hei en dass egal wie das Speicher abbild aussieht Feld Matrix oder Volumen die Threaddimensionierung dem angepasst werden muss Dann kann ein Thread ohne komplexere Adressberechnung sein Ele ment aus dem Speicher bearbeiten Als Beispiel wieder die Kernelfunktion vom Quell textabschnitt 2 1 in diesem Fall existieren genau so viele Threads wie Matrixelemente die in der gleichen Struktur angelegt sind Bei einem Speicherzugriff auf das nullte Ele men
90. en werden alle ste nummeriert und als Transferteil im Bild mit gr n markiert Zuletzt muss begin nend bei den InstTransfer store Operation dessen Transferteil r ckw rts durchlaufen wer den Wenn eine Operation gefunden wird die bereits der Pipeline zugeordnet wurde bricht die weitere Suche auf diesem Zweig ab analyseMemAccessPattern ist ein Platzhalter f r die n chste Version um die statischen bzw variablen Zugriffsmuster auf dem Speicher zu analysieren Im Prototyp wird diese Methode nicht ben tigt da die Anforderung existiert dass ausschlie lich statische Speicherzugriffe im Quelltext erlaubt sind 5 3 5 Parsebaum bersetzung Dieser Abschnitt entspricht dem bersetzungsschritt 2 aus dem Blockschaubild 5 4 5 3 5 1 SSA AST zu Block AST Nach der Analyse kann eine Instanz der Klasse sourcecompiler angelegt werden siehe Bild 5 10 Die ffentliche Methode compile generiert aus dem SSA AST einen neuen Block AST der aus Knoten besteht die der VHDL Pipeline entsprechen Dabei wird der SSA AST Knoten f r Knoten durchlaufen um den neuen Baum entstehen zu lassen SourceCompiler vhdlAst latencyLayerArray insertBlockAst layerDetermineRec lt VhdIBlock gt SourceAnalyser latencyFillLatencyArray latencySetinputDelay Bild 5 10 Auszug Klassendiagramm um die Compilerklasse Mit einer bersetzungstabelle siehe n chster Abschnitt wird auf jeden Knoten des SSA AST
91. enenverteiler engl crossbar switch Scheduler global geteilter Speicher Bild 3 6 Schematische Architektur der Speicher und der prozessierenden Elemente die Threadprozessoren aus einfachen fiinfstufigen MIPS Multithread Prozessoren mit dynami schen Prozess Schedulern implementiert worden Jeder Prozessor ist parametrisiert so dass die Anzahl der Threadprozessoren und die Bitbreite je nach FPGA Gr e gew hlt werden kann Weiter haben die Autoren die Idee spezifische Instruktionseinheiten f r eine Leistungssteige rung hinzuzuf gen bersetzungsvorgang Um den Programmieraufwand einer bersetzung gering zu halten werden so viele existierende Kompiliertechiken und Programmteile verwendet wie nur m g lich Zum einen spart dies viel Entwicklungszeit und zum anderen erh ht dies die Zuverl ssig keit da die Kompilierer Programmteile bereits weitl ufig genutzt werden Der bersetzungs vorgang ist in Bild 3 7 dargestellt Es gibt zwei Quellen f r die bersetzung einmal die Open CL Kernelfunktionen und einmal die Laufzeitumgebung die in unterschiedlichen Sprachen geschrieben werden k nnen Das GCC Frontend kann beide Quellen bersetzen W hrend die Laufzeitumgebung die LLVM Optimierer durchlaufen und effiziente bersetzte Objekte gene rieren werden die bersetzten Kernelfunktionen einer statischen Speicherzugriffsanalyse un terzogen Somit kann der Datenfluss zur Kompilierzeit optimal konfiguriert werden wodurch die O
92. entierungen Der Test PC auf dem die Messung durchgef hrt wurde besitzt ein Intel Core 2 Quad CPU Q6600 mit 2 4 GHz Taktrate 4 GBytes DDR2 Speicher mit 1066MHz Speichertakt Die Leistung der Grafikkarten werden im n chsten Abschnitt vorgestellt 111 6 Ergebnisse und Diskussion 6 1 2 Skalierung des Algorithmusses Die theoretischen Spitzenleistungen der drei Grafikkarten und der CPU sowie die erreichten Effizienzen der Haralick Bildmerkmalsimplementierung sind in Tabelle 6 2 aufgelistet Effizienz Effizienz Rechenleistung Speicherdurchsatz Rechnen Speicher GFLOPS GBytes s leistung durchsatz 1 CPU Q6600 theor Maximal 9 6 17 Implementierung 0 18 1 9 2 GPU 8800GTX theor Maximal 345 86 Implementierung 3 36 10 6 0 9 12 4 3 GPU GTX280 theor Maximal 622 141 Implementierung 8 20 8 1 3 14 8 4 GPU GTX480 theor Maximal 1345 177 Implementierung 15 37 7 1 1 21 3 Tabelle 6 2 Leistungsvergleich der Architektur mit der Implementierung Die Spalte Rechenleistung zeigt die theoretisch maximale Rechenleistung der jeweiligen Ar chitektur und die gemessenen Werte der Implementierung Ebenso wurde der maximale und der gemessene Speicherdurchsatz angegeben Die Messungen sind Durchschnittswerte f r den gesamten Algorithmus inklusive des Datentransfers zwischen GPU und CPU D h es gibt ein zelne Kernelfunktionen die sehr effizient rechnen bzw effiziente Sp
93. epr sentation von LLVM als AST verwenden k nnen Mit einem LLVM Programm l sst sich der LLVM AST aus dem LLVM IR SSA Quelltext generie ren Einen Parser f r LLVM IR SSA zu entwickeln ist einfach weil die assembler hn lichen Instruktionen zeilenweise zu lesen und zu interpretieren sind Den komplexeren LLVM AST zu verwenden weil er schon existiert war nicht Grund genug f r dessen Verwendung Weiter kann der eigene Parser gezielt f r die VHDL bersetzung entwi ckelt werden Eine nderung des LLVM AST f r VHDL w re weniger einfach gewesen 2 Aus dem SSA AST wird ein weiterer Syntaxbaum der Block AST generiert Die Kno ten des neuen Syntaxbaums sind keine Instruktionen mehr sondern Bl cke aus einem Baukastensystem f r Hardware Pipelines F r einige Instruktionen ist die Abbildung auf die Bl cke direkt m glich Hierbei zahlt es sich aus auf SSA Quelltext gesetzt zu ha ben da dieser einen unendlicher Registervorrat besitzt vergleichbar mit Signalen im 70 5 3 VHDL Kompilierer VHDL Quelltext Jedes Ergebnis wird in einem eindeutigen Register hinterlegt 35 add i32 30 33 entsprechend gibt es f r jedes Ergebnis in VHDL ein eindeutiges Signal s5 lt sO s3 3 Die Pipeline enth lt weitere Logik die neben dem Block AST existieren muss In diesem Schritt werden die Latenzzeiten berechnet und eventuelle Verz gerungsglieder hinzuge f gt Ebenso m ssen die Verbindungsglieder zum Rahmendesign konfiguriert
94. er CPUs im Bereich der parallelen Algorithmen erwartet wie z B bei der Bildverarbeitung Letzteres ist der kos teng nstigere und energiesparsamere Ansatz Eine Erweiterung des Rechen Clusters mit mehr Knoten und Beschleunigerkarten wird dadurch nicht ausgeschlossen 1 4 Forschungsfragen In dieser Arbeit werden zwei Themengebiete behandelt Erstens die Beschleunigung des Ha ralick Bildmerkmalalgorithmuses um die Geschwindigkeitsdiskrepanz zwischen der Daten aufnahme und Offline Prozessierung zu reduzieren Zweitens die Entwicklung einer verein fachten Programmierung der FPGA Karte f r die Online Prozessierung der aufgenommenen Bilddaten Mit einzelnen Fragen werden zentrale Forschungsfragen rund um das jeweilige Themengebiet entwickelt die in dieser Dissertation beantwortet werden e L sst sich der Haralick Bildmerkmalalgorithmus parallelisieren e Welche Co Prozessoren w rden f r eine Beschleunigung in Frage kommen Welche Re chenarchitektur passt besonders gut f r die Beschleunigung 1 Einf hrung und Ziele des Viroquant Projekts e Welche Beschleunigerkarte ist besser geeignet Grafikkarte oder FPGA Karte e L sst sich der Algorithmus im Bereich von zwei bis drei Gr enordnungen gegen ber einer CPU beschleunigen um den Biologen Wartezeiten zu ersparen e Welche Einflussfaktoren behindern eine weitere Beschleunigung e Rechtfertigt ein weiterer Arbeitsaufwand eine weitere Beschleunigung e Wie verh lt s
95. ereich von zwei bis drei Gr enordnungen Folgender Abschnitt zeigt einen Ansatz der den Anforderungen gerecht werden kann 3 2 2 Fazit f r eine Beschleunigung Diese Arbeit strebt den Ansatz an mit GPUs den Haralick Algorithmus inklusive aller 13 Bildmerkmale zu beschleunigen Folgende Gr nde sprechen daf r e Der Algorithmus enth lt berwiegend Flie kommaberechnungen die auf der GPU ein fach zu rechnen sind e Die Entwicklungszeit den Algorithmus auf die GPU zu parallelisieren ist k rzer als die Entwicklungszeit eines FPGA Designs auf der neusten Hardwaregeneration mit den fehlenden Bildmerkmalen e Die GPU bersteigt die theoretische maximale Rechenleistung der Flie kommaberech nungen eines FPGAs e Wegen der Offline Prozessierung ist es nicht wichtig die Berechnung nahe an die Sensor Elektronik zu kn pfen Auch der Energiebedarf der GPU ist im Rechencluster nicht kri tisch Eine h here Latenzzeit bei der blockweisen Berechnung ist bei der Offline Prozes sierung ebenfalls nicht relevant 3 3 Kompiliererentwicklung Die Online Prozessierung hat die Anforderung den FPGA Algorithmus den Bed rfnissen der Biologen entsprechend ver ndern zu k nnen Hardwarebeschreibungssprachen mit denen man den Algorithmus ver ndern kann sind f r die Biologen zu komplex Aus diesem Grund muss ihnen eine einfache Sprache zur Verf gen gestellt werden die sich in ein Hardwaredesign ber setzen l sst Dieser Abschn
96. ert werden damit ein Programm unterschiedli che OpenCL Ger te z B GPU und FPGA verwenden kann A OpenCL FPGA Beispielanwendung 1 void TestProgramKernelExecVecAdd const unsigned charx fileContent const size_t N o o o na wu v JO Ua A UDN 20 21 22 23 24 25 26 fileSize bool binary el_int status CL_SUCCESS cl_uint number 0 cl_device_id device NULL cl_context context NULL cl_program program NULL cl_kernel kernel NULL cl_command_queue queue NULL cl_mem meml NULL cl_mem mem2 NULL unsigned int h_meml NULL unsigned intx h_mem2 NULL unsigned int i 0 unsigned int Z 0 status clGetDeviceIDs NULL CL_DEVICE_TYPE_ACCELERATOR 1 amp device amp number if status CL_SUCCESS coutAndExit Could not find a device within platform status context clCreateContextFromType NULL CL_DEVICE_TYPE_ACCELERATOR NULL NULL amp status if status CL_SUCCESS coutAndExit Could not create context from Type status if binary program clCreateProgramWithBinary context 1 amp device amp fileSize amp fileContent NULL amp status if status CL_SUCCESS coutAndExit Could not create program with binary status else program clCreateProgramWithSource context 1 const char x 8fileContent amp fileSize amp status if status CL_SUCCESS coutAndExit Could not create program
97. erte und es entsteht mit dem Datentyp float eine Matrix mit der Gr e von 4096 x 4096 x 4Bytes 64M Bytes Speicherbedarf In einen 1024M Bytes gro en Grafikkartenspeicher w rden nur 16 Co Matrizen hineinpassen was wiederum die GPU nur zu einem Bruchteil auslasten w rde F r eine mas sive Parallelisierung des Algorithmusses m ssen die Co Matrizen kleiner werden Tats chlich sind die Co Matrizen nur sp rlich besetzt Das liegt daran dass die Zellbilder nicht rein zuf llig sind und die Pixelpaare bevorzugte Grauwerte haben Beispielsweise hat der Zel lenrand nur einen kleinen Wertebereich was bedeutet dass der Zellenrand rechts und links sehr hnlich ist Das gleiche trifft f r den Zellkern zu Besonders der Hintergrund des seg 49 4 Haralick Algorithmus GPU beschleunigt mentierten Zellbildes hat berall den gleichen Intensit tswert Allgemein entstehen bei der Co Matrix Berechnung mehr oder weniger bevorzugte Regionen in denen die Pixelpaare gez hlt werden w hrend andere Regionen komplett leer sind Bild 4 1 a zeigt eine sp rlich besetzte Co Matrix 0 4095 a b Bild 4 1 Bin rbild einer vollen Co Matrizen a und einer gepackten b Wei e Pixel entspre chen dem Wert Null schwarze Pixel einem von Null verschiedenen Wert Um Speicherplatz einzusparen werden alle Zeilen Aufgrund der Symmetrie auch Spalten die ausschlie lich Nullwerte enthalten entfernt Bild 4 1 b zeig
98. estandteile f r einen OpenCL FPGA bersetzer 66 5 3 VHDL Kompilierer 1 Die OpenCL Laufzeitumgebung ist eine Bibliothek mit Funktionen aus dem OpenCL Standard 36 siehe Bild 2 4 aus dem Grundlagenkapitel In der Bibliothek ist ein Kom pilierer integriert der Kernelfunktionen in unserem Fall f r den FPGA bersetzt All gemein verwaltet die Laufzeitumgebung die bersetzung und die Ausf hrung der Ker nelfunktionen die Speichertransfers und die Registerzugriffe auf den FPGA 2 Der VHDL Kompilierer besteht aus drei Programmen Clang Frontend LLVM IR Optimierer und VHDL Backend die zusammen eine bersetzungskette bilden In ihr wird eine OpenCL Kernelfunktion in eine VHDL Pipeline bersetzt 3 Das Rahmendesign bietet f r die generierte VHDL Pipeline einen Rahmen mit umlie gender Logik Speicherkontroller PCIe Core mit DMA Engine und Kontrolllogik Erst mit ihr kann die Pipeline mit Daten gef llt und eine Berechnung gestartet werden Der n chste Abschnitt 5 3 beginnt mit dem VHDL Kompilierer und beschreibt die Glieder der bersetzungskette sowie deren Implementierung Im anschlie enden Abschnitt 5 4 wird gezeigt welche Logik im Rahmendesign f r eine Ausf hrung notwendig ist und welche Im plementierungsans tze gew hlt worden sind Erst im letzten Abschnitt 5 5 nachdem das Rah mendesign bekannt ist wird die Implementierung der Softwareschnittstellen mit den OpenCL Funktionen beschrieben 5 3 VHDL Ko
99. f hrung ben tigen Handel C und CHiMPS oder ein Algorithmus versucht den seriellen Quelltext zu parallelisieren Trident Die paral lelen Sprachen haben den Vorteil bereits ein Paradigma zu liefern wie die Parallelit t auf dem FPGA abgebildet werden kann e F r den Biologen w rden die parallelen Sprachen OpenCL bzw CUDA und die Micro soft Accelerator Bibliothek in Frage kommen Sie sind leicht zu erlernen und erfordern nicht zwingenderma en Hardwarekenntnisse F r alle Genannten gibt es eine FPGA Entwicklung Das Derivat FCUDA CUDA f r den FPGA ben tigt allerdings zus tzli che Anweisungen f r den bersetzungsprozess und scheidet deswegen aus OpenRCL OpenCL f r rekonfigurierbare Logik hat alle w nschenswerten Eigenschaften mit einer Ausnahme n chster Punkt Der Microsoft Accelerator ist nicht so bekannt und verbrei tet wie OpenCL oder CUDA und die Accelerator Bibliothek ist im Vergleich nicht intuitiv zu erlernen OpenCL ist von der Programmierwelt besser akzeptiert e Die OpenRCL und die FCUDA Entwicklung verfolgt den Ansatz einer Rechnerarchi tektur wie die einer GPU vielen Threadprozessoren nachzuahmen Die Berechnungen werden auf mehrere Threadprozessoren verteilt wobei jede Operation mehrere Takte Re chenzeit beanspruchen kann FPGAs entfalten ihr Potenzial mit einer Hardwarepipeline in der alle Operationen parallel in jedem Takt berechnet werden Auch die Latenzzeit ist in einer Pipelineverarbeitung ge
100. g geachtet um Funktionalit t hinzuf gen zu k nnen Folgende Vorschl ge w rden den OpenCL FPGA Kompilierer verbessern e Die Verwendung eines OpenCL Frontends w rde die Nutzung des Shared Speichers er m glichen wof r es ein Konzept im FPGA Design zu entwickeln gilt e Das Hinzuf gen weiterer Sprachelemente wie if e1se Strukturen und Schleifen w rde die bersetzung komplexerer Programme erm glichen e Die Erweiterung der Speicherzugriffsanalyse und der entsprechenden Adresskomponen te l st die Einschr nkung bestimmter Speicherzugriffsmuster 121 7 Fazit und Ausblick 122 Zur Unterst tzung von loat Datentypen muss eine Flie kommabibliothek dem Kompi lierer und dem VHDL Block Vorrat hinzugef gt werden Komplexere Kernelfunktionen erh hen den Ressourcenbedarf der Pipelineregion auf dem FPGA der vergr ert werden muss Sofern die FPGA Ressourcen ausreichen k nnte die Pipeline mehrfach implementiert werden um unterschiedliche work groups parallel abarbeiten zu lassen was den Re chendurchsatz erh hen w rde Das Rahmendesign k nnte angepasst werden um gleichzeitig Rechnungen und Daten transfer zwischen Host und OpenCL Ger t zuzulassen Die vereinfachte Implementierung der Laufzeitumgebung m sste erweitert werden um dem vollen OpenCL Standard zu entsprechen Es gibt einen ICP Treiber um unterschiedliche OpenCL Laufzeitumgebungen nebenein ander anzubieten Sie m sste implementi
101. g wird f r jede generierte Co Matrix ein eigener Block geschaffen Jede Operation die auf die Bl cke angewendet wird impliziert die parallele Ausf hrung auf alle Co Matrizen In Bild 4 5 ist die Struktur der GPU Version in einem Struktogramm erl utert Die u erste Schleife iteriert ber alle Zellen des Multizellbildes Der Unterschied zur CPU Version ist dass mit jedem Scheifendurchlauf C Zellen gleichzeitig gelesen werden Innerhalb der Schleife werden f r alle C Zellen wiederum gleichzeitig alle Co Matrizen AD erstellt F r jede Sichtweise des Winkel A und der Distanz D wird eine eigene Co Matrix generiert Aus diesem bel uft sich die Anzahl an Kombinationen zu Ax D AD 56 4 3 GPU Implementierung Die Anzahl der CUDA Bl cke bel uft sich auf Cx AD in denen die Zwischenergebnisse und Bildmerkmale berechnet werden Am Ende des Schleifenk rpers werden die Bildmerkmale gespeichert und die Co Matrizen gel scht Schleife ber alle Zellen C in parallel generiere Matrizen C AD in parallel Berechne alle Zwischenergebnisse und Bildmerkmale C AD in parallel speichere Bildmerkmale l sche Matrizen Bild 4 5 Struktogramm der parallelen GPU Version Die Berechnung des Haralick Algorithmusses mit allen Bildmerkmalen und Zwischenergeb nissen ist Komplex Implementiert in eine einzige Kernelfunktion w rde das den Quellcode seitenlang werden lassen und die ben tigten Ressourcen der GPU w rden nicht ausrei
102. gbare OpenCL Implementierungen auf Datum Entwickler Architektur Quelle 20 04 2009 NVidia GPU 11 05 08 2009 AMD GPU und CPU 9 30 06 2010 IBM Power und Cell 12 13 09 2010 Intel CPU 13 Tabelle 3 2 OpenCL Implementierungen aufgelistet nach Erscheinungsdatum Apple 10 entwickelte f r die GPUs von Nvidia und AMD eine eigene OpenCL Implementie rung die im Betriebssystem Mac OS X Snow Leopard verankert ist ARM S3 und VIA bieten bereits Produkte an die OpenCL f hig sind jedoch existieren noch keine Softwarepa kete f r potenzielle Entwickler Eine vollst ndige Liste von OpenCL Implementierungen bzw Produkten die OpenCL f hig sind gibt es auf der Seite von Khronos 14 OpenCL f r den FPGA nutzbar zu machen wie es das OpenRCL Projekt 3 4 3 2 gemacht hat 43 1 2 3 4 5 6 7 Aa Bw N Oo amp VD o N 0 3 Stand der Technik bietet den Vorteil keine weitere Programmiersprache lernen und keine weitere Entwicklungs zeit f r eine Portierung investieren zu m ssen 3 4 4 2 Microsoft Accelerator Die Ver ffentlichung Computer without Processors 70 vertritt die Meinung dass der Re chenbedarf von Cloud Computing am besten mit einer Mischung vieler Rechenarchitekturen CPU GPU und FPGA erf llt wird und dass in der Zukunft h ufiger Mischarchitekturen zum Einsatz kommen Das Entwicklungswerkzeug Accelerator 72 von Microsoft zielt auf die Unabh ngigkeit der Ar
103. h der Pipeline 6 Optional kann die acht Bit gro e Pipeline Modul ID mit devsetprid aus dem Kontroll register ausgelesen werden um den Austausch zu verifizieren 110 6 Ergebnisse und Diskussion Nachdem die Implementierungen der GPU Beschleunigung offline Prozessierung im vierten Kapitel dargelegt und die Implementierung des OpenCL Kompilierers online Prozessierung im f nften Kapitel erl utert wurde werden in diesem Kapitel die Ergebnisse zu den Implemen tierungen zusammengetragen 6 1 GPU Beschleunigung 6 1 1 Geschwindigkeitsgewinn Folgende Tabelle 6 1 zeigt die Ausf hrungszeiten des Haralick Texturen Algorithmus die al le das gleiche Multizellbild f r ihre Berechnung verwendet haben Verglichen wird die Ori ginalsoftware mit einer optimierten Software Version einer ersten GPU Version I und einer optimierten GPU Version II Mit enthalten sind die Ausf hrungszeiten dreier unterschiedlicher Grafikkarten die im Zeitraum dieser Arbeit erschienen sind Ausf hrungs Faktor zu Faktor zu Faktor zu Zeit s 1 2 3 1 Original Software Version 2378 2 Optimierte Software Version 214 11x 3 GPU Version I 8800 GTX 11 1 214x 19x 4 GPU Version I GTX 280 6 6 360x 32x 1 7x 5 GPU Version II GTX 280 4 65 511x 46x 2 4x 6 GPU Version II GTX 480 2 55 933x 84x 4 3x Tabelle 6 1 Ausf hrungszeiten und Beschleunigungsfaktoren im Vergleich aller vorgestellten Implem
104. hen In Hardware wurde dies mit Multiplexer und Demultiplexer realisiert kontrolliert vom Transfermodus Im DMA Transfermodus be kommt die PCle Einheit Signale mit dma Pr fix den exklusiven Zugriff auf die Spei cherschnittstelle Die Speicherzugriffe werden direkt an die Speicherschnittstelle wei tergeleitet Im BUS Transfermodus erhalten die Transferbl cke Signale mit bus Pr fix die Kontrolle ber den Speicher Aufgrund der Einfachheit der Implementierung wur de darauf verzichtet beiden Schnittstellen gleichzeitig Speicherzugriffe zu erm glichen D h der Host kann erst auf den Speicher zugreifen wenn die Pipeline nicht rechnet Am Lese Bus und am Schreib Bus k nnen mehrere Transferbl cke die Slaves ange schlossen sein Die Anzahl der an den Bussen angeschlossenen Slaves wird von den Signalen bus_def_rd_size und bus_def_wr_size aus der Pipeline von der zus tzlichen Logik geliefert In der Speicherverwaltungseinheit sitzt der Bus Master der ber die sel Signale alle Slaves nacheinander selektiert die ihrerseits mit den done Signalen ent weder einen Speicherzugriff anfordern oder mitteilen dass sie keinen Bedarf eines Spei cherzugriffs haben Der Bus Master selektiert alle Slaves abwechselnd auf dem Lese Bus und dann auf dem Schreib Bus da es keine Notwendigkeit des Full Duplex Betriebs gibt obwohl die M glichkeit mit zwei getrennten Bussen best nde Wenn ein Slave des Schreib Busses selektiert wird kann dieser mit de
105. hreib Schnittstelle Sie besteht aus Datenleitungen Maskierungsleitungen und Kon trollleitungen Die minimale Datenmenge von 512 Bit muss immer in zwei aufeinander folgenden Takten mit 256 Bit Datenworten gesendet werden Ebenso kann die Schnitt stelle mitteilen dass sie gerade keine Daten akzeptiert beispielsweise wenn der interne Schreibpuffer voll ist Lese Schnittstelle Sie besteht aus den Datenleitungen und Kontrollleitungen Die ange forderten Daten werden ebenso wie beim Beschreiben als zwei aufeinander folgende 256 Bit Worte geliefert F r den Zweck dieser Arbeit besitzt die Benutzerschnittstelle des Speicherkontroller mehrere Nachteile Die Benutzerschnittstelle arbeitet mit fest eingestellten 200MHz Die Kommandoschnittstelle und die Schreibschnittstelle m ssen synchronisiert ange sprochen werden Es existiert ein spezifisches Zeitverhalten beim Schreiben was einen Schreibvorgang kompliziert gestaltet und eine Zustandsmaschine bedarf Beim Lesen einer Adresse mit Byteoffset werden die Daten innerhalb der 256 Bit Worte umsortiert was eventuell gar nicht gew nscht ist Die selben Adressleitungen m ssen f r das Lesen und das Schreiben verwendet werden Eine zwischengeschaltete Zugriffs Komponente besitzt vereinfachte Schnittstellen logisch in zwei Teile aufgeteilt eine zum Lesen und eine zum Schreiben Beide habe ihre eigenen Adress leitungen und Wartesignale die einer FIFO Schnittstelle hneln Die
106. hung zu Bild 5 5 Klassendiagramm des VHDL Backends 72 5 3 VHDL Kompilierer In den folgenden Abschnitten wird die Implementierung der bersetzungsschritte 1 bis 4 er l utert und anhand von Beispielen verdeutlicht 5 3 3 Parsebaum Generierung Dieser Abschnitt entspricht dem bersetzungsschritt 1 aus dem Blockschaubild 5 4 Bild 5 6 zeigt die Klasse sourceParser mit allen n tigen Klassen um den SSA AST zu generieren Das Parsen geschieht in der Methode parse in einer Schleife die jede Zeile des SSA Quelltextes mit drei Hilfsmethoden auswertet SourceAttrib InstructionType Bild 5 6 Klassendiagramm des Parsers mit unvollst ndiger Bezeichnung der Instanzvariablen und der Methoden findInstructionType identifiziert die Instruktion der Quelltextzeile und ordnet sie einer Ka tegorienummer zu In der Klasse instructionset sind alle Instruktionen als string gespeichert in der die Methode nach dem entsprechenden Eintrag suchen kann Die Klasse InstructionType enth lt die Zuordnung der Kategorienummer eine Art Enumerationstyp mit der sich die In struktion besser identifizieren und vergleichen l sst als mit einem string parseInstruction erstellt anhand der Kategorienummer eine spezialisierte Instanz von der abstrakten Klasse Instruction im Bild nicht dargestellt Das Parsen also das Auswerten der Quelltextzeile passiert im Konstruktor der abgeleiteten Klasse Jede Instruktion h
107. i gen sie f r ihre Ausf hrung Ein Beispiel einer OpenCL Anwendung wie die Funktionen der Laufzeitumgebung verwendet werden befindet sich im Anhang A 2 2 Rekonfigurierbare Hardware 2 2 1 Aufbau eines FPGAs Die Abk rzung FPGA bedeutet field programmable gate array frei bersetzt hei t das feld programmierbarer Logikschaltkreis Das Bild 2 5 zeigt den schematischen Aufbau eines FPGAs Er besteht aus vielen einzelnen Logikzellen LZ Blockspeichern Block RAM Multiplizie rern Ein und Ausgabebl cken IOB und einem programmierbaren Verdrahtungsnetzwerk Wie ein FPGA aufgebaut wird ist im Nachschlagewerk Das FPGA Kochbuch 78 gut er kl rt 17 2 Grundlagen BEN device_id e ee Ti ane context cl al coun Pe aa i clGetDevicelDs gt clCreateContextFromType gt lt W u clCreateCommandQueue cLeommand_queve command cLeommand_queve cl_kernel clSetKernelArg clEnqueueWriteBuffer ee 4 clEnqueueNDRangeKernel clEnqueueReadBuffer Bild 2 4 Verwendete Strukturen und Funktionen eines minimalen OpenCL Programms ein schlie lich einer bersetzung mit Kernelaufruf Pfeile sind als generierendes Ergeb nis zu interpretieren Punkte als Abh ngigkeit Hi ate uta HHH mi TT CHIH irh Eile ptr C ant a db ele an m me tuts tuts un E mE Eu N zus
108. ich die Entwicklungszeit einer GPU L sung zu einer FPGA L sung e Welche Programmiersprachen k nnen genutzt werden um die Entwicklungszeit eines FPGA Designs zu beschleunigen und zu vereinfachen e Welchen Vorteil besitzt eine parallele Programmiersprache wie z B OpenCL gegen ber einer seriellen Programmiersprache zur FPGA Beschreibung e Worin liegen die Unterschiede einer Architektur in der eine Pipeline entsteht gegen ber einer Architektur mit gleich bleibendem Rechenwerk e Wie sieht die Struktur eines Kompilierers f r eine Pipelinegenerierung aus e Welche weitere Logik wird an den Schnittstellen zur generierten Pipeline ben tigt F r das Themengebiet der Onlineprozessierung ergibt sich folgende zentrale Forschungsfra ge Wie sehen die Bausteine einer bersetzungskette aus die von der parallelen Pro grammiersprache OpenCL in eine VHDL Hardwarepipeline f r FPGAs bersetzen und welche Vorteile bietet diese Struktur Zusammenfassend l sst sich f r die Offlineprozes sierung fragen Wie l sst sich der Bildmerkmalalgorithmus von Haralick auf einer GPU beschleunigen welcher Beschleunigungsfaktor gegen ber einer CPU ist zu erreichen und welche Einflussgr en tragen ma geblich zur Beschleunigung bei 2 Grundlagen 2 1 Grafikkarten als Rechenbeschleuniger 2 1 1 Geschichtliche Entwicklung Die Graphics Processing Unit GPU ist ein Prozessor der die Central Processing Unit CPU f r grafisch
109. ickeln 5 4 2 PCle Core und DMA Engine Mit dem PCIe Core von Xilinx 8 1 wird ein Kommunikationskanal zwischen der Host Software und dem Hardware Design ber den PCI Express Bus hergestellt Mit dem Coregenerator von Xilinx wird der vorhandene Hardware PCle Core auf dem Virtex6 Chip konfiguriert Die we sentlichen Eigenschaften der Konfiguration sind e Der PCIe Core wird in der Xilinx Version 1 3 verwendet Aktuellere Versionen laufen nicht auf der verwendeten Entwicklungssteckkarte da dessen FPGA ein Prototyp Engi neering Sample ist und in diesem Fall eine eingeschr nkte Funktionalit t besitzt 91 5 OpenCL zu FPGA bersetzer e Es werden vier von m glichen acht Lanes verwendet e F r die Bus Kommunikation wird der PCle Standard der Generation 1 1 zugrunde ge legt Jede Lane hat eine maximale Ubertragungsrate von 250MB s 1000MB s entspre chend f r vier Lanes e Die Schnittstelle des PCle Cores soll mit einer Frequenz von 100 MHz angesprochen werden e Der PClIe Core ben tigt eine Lizenz die der Universit t geh rt Wenn die CPU Daten ber den PCIe Bus sendet oder liest PIO Mode kann eine Datenrate in der Gr enordnung von zehn Megabytes pro Sekunde erwartet werden Aus diesem Grund wird eine DMA Engine an den PCle Core geheftet die Datenraten von bis zu 790MB s lesend und 507MB s schreibend unterst tzt Der Unterschied zur theoretisch maximalen Bandbrei te besteht in den notwendigerweise zu ber
110. ie Harder Apichat Suratanee Karl Rohr Rainer Koenig und Reinhard Maenner Haralick s Texture Features Computation Accelerated by GPUs for Biological Applications In To appear in Modeling Simulation and Opti mization of Complex Processes Proceedings of the Fourth International Conference on High Performance Scientific Computing March 2 6 2009 Hanoi Vietnam Seite 127 138 Springer Verlag Berlin Heidelberg 2011 doi 10 1007 978 3 642 25707 0_11 Khronos Group OpenCL Quick Reference Card URL http www khronos org files opencl quick reference card pdf 2009 Khronos OpenCL Working Group The OpenCL Specification version 1 0 48 6 October 2009 R M Haralick Statistical and structural approaches to texture Proceedings of the IEEE Volume 67 5 786 804 1979 doi 10 1109 PROC 1979 11328 R M Haralick und K Shanmugam Computer Classification of Reservoir Sandstones Geoscience Electronics IEEE Transactions on Volume 11 4 171 177 1973 doi 10 1109 TGE 1973 294312 Robert M Haralick K Shanmugam und Its Hak Dinstein Textural Features for Image Classification Systems Man and Cybernetics IEEE Transactions on Volume 3 6 610 621 1973 John Hennessy und David Patterson Computer Architecture A Quantitative Approach Morgan Kaufmann 2007 Allen Holub Compiler design in C Prentice Hall Inc Upper Saddle River NJ USA 1990 Agility Design Solutions Inc Handel C Language Reference Manua
111. ie entsprechende Stelle geschoben werden Das ist eine gute L sung um 44 3 4 Software Hardware Kompilierer Speicherzugriffe zu reduzieren da die Daten einmal aus dem Speicher gelesen werden aber aus dem Schieberegister mehrfach verwendet werden Die Berechnung wird in Zeile 15 auf dem Zielger t zur Ausf hrung gebracht je nachdem welches Target instantiiert wurde Es gibt Implementierungen f r FPGAs Fpcatarget f r GPUs px9Target die auf die Shadersprache DirectX Version 9 aufsetzt und CPUs x64MulticoreTarget f r Vielkern CPUs mit Ausnut zung der SSE3 Vektorisierung Die Ver ffentlichung demonstriert in den Ergebnissen wie die Beschleunigung der DX9 Ver sion auf Ati und NVidia GPUs mit l nger werdenden Datenreihen linear skaliert W hrend die Beschleunigung mit der Vielkern CPU Implementierung aufgrund der Cachegr e ein Maxi mum hat und die Beschleunigung mit l nger werdenden Datenreihen mit kleinerem Beschleu nigungsfaktor nicht weiter ansteigt Ziel der Entwicklung ist es Algorithmen in einer Sprache zu schreiben und f r viele Zielger te nutzen zu k nnen Entwicklungszeit bzw Portierungszeit einzusparen M chte man das Maximum an Leistung des Zielger tes aussch pfen bleibt nichts anderes brig als die Nativ Sprachen CUDA VHDL der Zielger te zu benutzen 3 4 5 Fazit der FPGA Sprachen Diese Arbeit strebt den Ansatz an einen FPGA f r die Online Prozessierung einzusetzen Fol gende Gr n
112. ie sich aus der Basisadresse und dem Adressoffset ergeben Der Adressoffset wird aus den Parallelindizes hergeleitet PARAM_POSITION TRANS_MODE LINEAR_GLOBAL clkUser work item id rst idx_local_size_ x y z Adress berechnung idx_work_group_size_ x y z idx_work_group_id_ x y z base_pos base_rdy acalc_addr acalc_request req acalc_rdy acalc_update UP acalc_burst acalc_next acalc_valid acalc_start_grou group acalc_fin Bild 5 14 Schematische Darstellung des VHDL Adressblocks Oben Generic Ports und unten die Entity Ports e Die Basisadresse ist erst zur Laufzeit bekannt und muss ber die base Schnittstelle mit geteilt werden Der Genericparameter gibt der Komponente eine eindeutige Nummer zur Identifikation e Die Parallelindizes werden in aufsteigender Reihenfolge abgearbeitet In der Kompo 84 5 3 VHDL Kompilierer nente wird der kleinste Parallelindex inkrementiert und mit den brigen Indices der idx Schnittstelle verrechnet Zusammen mit dem Genericparameter der das Speicher zugriffsmuster bestimmt wird der Adressoffset bestimmt e Die Adressberechnung wird von einer Zustandsmaschine gesteuert Sie regelt auch die Kommunikation mit der acalc Schnittstelle nimmt Anfragen entgegen und gibt Be scheid wenn die Adressberechnung beendet ist Der funktionale Aufbau des Transferleseblocks ist in Bild 5 15 dargestellt Dieser Block hat die Aufgabe mit Hilfe des Adressblocks die D
113. ieder zu synchronisieren und am Ende die prozessierten Daten in den Grafikkartenspeicher zu kopieren In der Implementierung der GPU Version II werden in den meisten Kernelfunktionen die Operationen direkt auf dem Grafikkartenspeicher angewandt Lesen Ausf hren der Operationen und Speichern der prozessierten Daten geschieht ohne den Shared Speicher zu verwenden Synchronisation wurde hinf llig was zu einer beschleunigten Ausf hrung gef hrt hat Auch die Caches die in den neueren GPU Generationen vorhanden sind tragen dazu bei den Shared Speicher nicht mehr als Datenpuffer einsetzen zu m ssen Mit den gezeigten Ma nahmen reduzieren sich die Ausf hrungszeiten im Durchschnitt um den Faktor 1 4 Eine detaillierte Auflistung der Kernelfunktionen mit den Ausf hrungszeiten mit und ohne Optimierungen ist im Kapitel Ergebnisse 6 1 auf Seite 111 zu finden An dieser Stelle ist ebenso ein Geschwindigkeitsvergleich zur CPU und das Skalierungsverhalten mit der erreichten Rechenleistung in GFLOPS dargestellt 63 5 OpenCL zu FPGA bersetzer 5 1 Konzept Der beschriebene Ansatz f r einen OpenCL FPGA bersetzer aus Kapitel 3 4 5 wird hier ver feinert und gegen ber Alternativen abgewogen Einen FPGA einzusetzen ist notwendig weil eine FPGA Karte einen Datendurchsatz von mehreren hundert MBytes s aufnehmen kann Die Sensoren werden ber differentiale Signale LVDS mit dem FPGA verbunden Andere bliche Schnittstellen wie zum Beispiel USB o
114. ign Beide Instruktionen werden kombiniert und in InstSelect einen VHDL Block bersetzt InstComp BlockComp Wird mit einer Eins zu Eins Beziehung bersetzt Die SSA Instruktion parametrisiert die VHDL Komponente InstGetElemPtr Die Berechnung der Adresse passiert in der BlockAddr Komponente kombiniert mit der Ana lyse der parallelen Struktur InstTransfer BlockAddr Diese Komponente ist immer an eine BlockTransfer Komponente gekoppelt je nach dem welche Instruktion gemeint ist BlockTransRd Wird generiert wenn die Instruktion ein 1oad ist BlockTransWr Wird generiert wenn die Instruktion ein store ist Tabelle 5 2 bersetzungstabelle vom SSA AST zum Block AST Die OpenCL Kernelfunktion wird von vielen Threads durchlaufen genauso soll die Pipeli ne von vielen Datenelementen gestreamt werden Folglich existiert eine Verbindung zwischen dem n ten Datenelement mit dem n ten Thread aus dem work items Vorrat F r die Rechenpi 80 5 3 VHDL Kompilierer peline ist es nicht wichtig zu wissen von wie vielen Datenelementen Threads sie durchlaufen wird Aus diesem Grund wird die instca11 Klasse f r den Block AST nicht ben tigt Auch das Zugriffsmuster auf den Speicher ist f r die Rechenpipeline uninteressant Die Komponenten Block TransRd und BlockTransWr bieten lediglich einen lesenden bzw schreibenden Zugang zum Speicher In der BlockAddr Komponente wird der Index und die Zieladresse des aktu ell gestr
115. ilder l sst sich im allgemeinen sehr gut auf den Grafikkarten sowohl in der Pixel Ebene als auch in der Bild Ebene parallelisieren bzw beschleunigen Eine Anwendung f r die beschleunigte Offline Prozessierung ist der Haralick Bildmerkmalsalgorithmus der den Biologen gute Ergebnisse liefert daf r aber einen hohen Rechenaufwand besitzt Die GPU basierte Beschleunigung dieser Arbeit reduzierte die Re chenzeiten von Monaten auf Stunden Aufgrund des Skalierungsverhaltens der letzten GPU Generationswechsel versprechen zuk nftige Grafikkarten mit einem einfachen Austausch ei ne weitere Beschleunigung ohne den Algorithmus anpassen zu m ssen Die zu erwartende Beschleunigung mit der bereits erzielten ergeben keinen Bedarf den Algorithmus weiter ver bessern zu m ssen Eine weitere Beschleunigung w rde f r die Gesamtausf hrungszeit der Algorithmenkette aus der automatischen Bildanalyse keinen nennenswerten Vorteil bieten Eingesetzte CPU Version 1 Monat Optimierte CPU Version 3 Tage 8800 GTX GPU Version 2 Stunden GTX 480 GPU Version Il 45 Minuten Bild 7 1 Vergleich der Ausf hrugszeiten der unterschiedlichen Versionen Bild 7 1 illustriert die Ausf hrungszeiten des Haralick Algorithmuses mit unterschiedlichen Implementierungen Als Referenz dient die Originalversion die eingesetzt wurde um die Bild merkmale zu berechnen Sie ben tigt f r einen Datensatz einer bestimmten Gr e einen Monat 119
116. in ein Speicherabbild berf hrt und die Funktion berpr ft Die se Aufgaben werden von den drei Teilen erledigt Zeilenrekonstruktion Als vorprozessierenden Schritt werden Leerzeichen Tabulator Zeichen Zeilenende Zeichen und Kommentare gefiltert da sie f r die bersetzung in den meisten Spra chen keine Relevanz haben 21 2 Grundlagen Lexikalische Analyse Als erstes muss der Quelltext in Tokens zerlegt werden Tokens sind grammatikalische Bausteine z B Schl sselworte Sprachsymbole Bezeichner Operatoren usw Syntaktische Analyse Im zweiten Schritt werden die Tokens den Grammatikregeln der Spra che unterzogen und gepr ft ob deren Reihenfolge zul ssig ist Dabei wird ein Syntaxbaum AST abstract syntax tree im Speicher konstruiert Gibt es einen Regelversto bricht die syn taktische Analyse auch Parser genannt mit einer Syntax Fehlermeldung ab Semantische Analyse Zuletzt wird inhaltlich nach der Bedeutung und dem Sinn gepr ft Mit Hilfe von Attributen wird der Syntaxbaum mit zus tzlichen Informationen angereichert Varia blen und parametrisierte Funktionsreferenzen werden in einer Objektabelle aufgelistet und mit jeder Zuweisung einer Typpr fung unterzogen Widerspricht das Programm der Sprachlogik bricht der Kompilierprozess mit einer Semantik Fehlermeldung ab 2 3 2 Backend Das Backend ist der zweite Teil eines Kompilierungsprozesses der folgende Aufgaben um fasst Generierung ei
117. in passen Die Sch tzung ber cksichtigt ausschlie lich Integer Operatoren und eine durchschnittliche Nutzung aller existierender Pipelinebausteine Sofern eine Flie komma bibliothek in den Kompilierer einbezogen wird w re die Bausteindichte mit Flie kommaope ratoren wesentlich kleiner 6 2 5 Beispiel Applikationen Ein Multizellbild das vom Haralick Bildmerkmalsalgorithmus offline auf der GPU prozessiert wird muss einen planen Hintergrund der Bereich zwischen den Zellen mit Nullwerten haben Eine Variante den Hintergrund zu entfernen ist das Multizellbild 6 1 a mit einem segmen tierten Bin rbild 6 1 b zu multiplizieren Das Bin rbild hat den Hintergrundwert Null und innerhalb der Zellen den Wert Eins Die Aufgabe den Hintergrund zu entfernen l sst sich auf dem FPGA online prozessieren Fol gendes Quelltextbeispiel 6 1 zeigt die Kernelfunktion die dazu imstande ist und sich f r den FPGA bersetzen und ausf hren l sst __ kernel void removeBGwithSement Img intx multicell int segment int Z int w int idx get_global_id 0 116 OIDO ua a LODO a BF BW Ye 6 2 OpenCL Kompilier a Bild 6 1 Multizellbilder a verrauscht im Hintergrund und b bin r segmentiert Il int idy get_global_id 1 int index idy w idx Z index multicell index segment index Quelltext 6 1 Kernelfunktion die ein Multizellbild mit einem segmentierten Bin rbild multipl
118. in x oder exp x Speicher und Register Jeder TP hat unterschiedliche Speicherbereiche zur Verf gung Daten abzulegen Am schnellsten sind die 32786 Register mit je 4 Byte Speicher die auf alle TPs auf geteilt werden Der geteilte Speicher Shared Speicher ist f r den Datenaustausch zwischen unterschiedlichen TPs gedacht Es existiert ein 64 kBytes Shared Speicher der gleichzeitig als L1 Cache verwendet wird Weitere Details verschiedener Speicher sind im Abschnitt CUDA erl utert Die theoretisch maximale Rechenleistung berechnet sich durch 480 Thread prozessoren x 1 4 GHzx2 Operationen 1344 GFLOPS mit einfacher Genauigkeit SP single precision Die TPs sind mit 1 4GHz getaktet und mit der MAC Operation k nnen sie gleichzeitig zwei Ope rationen die Multiplikation und die Addition ausf hren Wenn in doppelter Genauigkeit DP double precision gerechnet wird ist die Rechenleistung um den Faktor acht kleiner 2 1 3 Programmiersprache CUDA Die Compute Unified Device Architecture CUDA ist eine Programmiersprache die auf C auf setzt Sie erm glicht das Programmieren von NVIDIAs Grafikprozessoren mit nicht grafischen Anwendungen CUDA besteht aus zwei Teilen den Kernel Funktionen die parallel tausende Threads auf der GPU ausf hren k nnen und der Laufzeitumgebung einem Funktionsumfang der Speichertransfers und Kernelaufrufe vom Host aus regelt um die GPU zu steuern In CUDA kann der Programmierer einen Algorithmus pa
119. ine Variable im Shared Speicher hinter legen Diese Einschr nkung hat Einfluss auf die F higkeiten des Kompilierers Daten k nnen zwischen den Threads nur ber den globalen Speicher ausgetauscht werden Urspr nglich war geplant ein OpenCL Frontend zu verwenden Zu Beginn der Entwicklung war bekannt dass OpenCL Entwickler daran arbeiten Clang f r OpenCL Kernelfunktionen zu erweitern Heute ist bekannt dass Apple und NVIDIA ein OpenCL Frontend basierend auf Clang und LLVM verwenden 17 um Quelltext in OpenCL f r ihre Plattformen zu berset zen Wir hatten darauf spekuliert dass das OpenCL Frontend mit der Zeit ver ffentlicht wird um es als Ersatz zu Clang verwenden zu k nnen was bisher nicht geschehen ist Clang in die ser Arbeit f r OpenCL Kernelfunktion abzu ndern war zeitlich keine Option denn der Fokus dieser Arbeit liegt auf dem neuartigen VHDL Backend Ein Vorteil LLVM IR zu verwenden ist es den umfangreichen Optimierer f r diese Zwischen sprache nutzen zu k nnen Bild 5 3 zeigt ein Beispiel wie die Kernelfunktion vom Clang Frontend in die Zwischensprache LLVM IR bersetzt wird und danach die Anzahl der Instruk tionen vom LLVM Optimierer reduziert wird Ohne Optimierer betr gt der LLVM IR Quell text ber 40 Zeilen mit lediglich 12 was einer Reduktion von 70 entspricht Die Kommando zeilen Dokumentation des Optimierers zeigt welche Verfahren f r die Quelltext Reduktion existieren Der Optimierer wird weg
120. iprozessor ausgef hrt durch andere CUDA Bl cke verdr ngt bzw beendet wenn alle Thread des Blocks das Ende der Kernelfunktion erreicht haben Die genaue Zuordnung der CUDA Bl cke zu den SMs und die Verteilung der Warps auf die Threadprozessoren passiert dynamisch zur Laufzeit und ist deswegen nicht vorhersagbar Die CUDA Bl cke werden in ihrer Ausf hrung verdr ngt wenn ihre Threads Speichertransfers initiieren und auf die Daten vom Speicherkontroller warten Da die Latenzzeiten viele hundert Taktzyklen dauern lohnt es sich w hrend dessen einen anderen CUDA Block auszuf hren Sofern ein Vielfaches an Bl cken zu den SMs existiert k nnen die Latenzzeiten der Speicherzugriffe hinter der Berechnungszeit anderer CUDA Bl cke vollst n dig versteckt werden Tabelle 2 1 zeigt alle verf gbaren Speicherbereiche die von den Kernelfunktionen genutzt wer den k nnen Register Die Programmierer k nnen auf die Register nicht direkt zugreifen Der CUDA Kompilierer verwaltet die Register und strebt an sie sehr effizient zu nutzen da sie die ge ringste Latenzzeit aufweisen Lokale Speicher Falls ein CUDA Block mehr Register ben tigt als vorhanden sind m ssen 11 2 Grundlagen Gr e Latenz Chip Cached GPU Zugriffs zeiten integriert Speicher Zugriff beschr nkung Register 8192 1 Ja Nein lesen schreiben pro TP Local global 1 800 Nein Ja lesen schreiben pro TP Shared 16 KB 2 Ja Nein lesen schreiben pro SM Global 1 5
121. it weniger Rechenoperationen ersetzt werden Die restlichen Bildmerkmale und Definitionen sind unver ndert und werden zur Vollst ndigkeit mit Namen aufgelistet Angular Second Moment 4 1 Contrast 4 2 Variance 4 4 Inver se Difference Moment 4 4 Sum Difference Average 4 6 Sum Variance 4 7 Sum Entropy 4 8 Entropy 4 9 Difference Variance 4 10 Difference Entropy 4 11 und Information Measurement II 4 13 Die meisten Bildmerkmale 4 1 4 4 4 6 4 8 und 4 10 4 13 h ngen von anderen Bild merkmalen und Zwischenergebnissen bzw der Definitionen ab Um eine zeitintensive doppelte Berechnung der Bildmerkmale und Zwischenergebnisse zu vermeiden muss die richtige Rei henfolge ermittelt werden Die Abh ngigkeiten wurden analysiert und in einem Graphen in Bild 4 3 dargelegt Jeder Kreis des Graphen ist ein Bildmerkmal f r dessen Berechnung alle anderen mit Linien verbundenen Kreise bzw K stchen mit Zwischenergebnissen berechnet sein m ssen da dessen Ergebnisse in die Berechnung mit einflie en Zum Beispiel muss erst das Ergebnis von Bildmerkmal 4 7 ermittelt werden weil Bildmerkmal 4 6 es f r seine Be rechnung ben tigt 53 4 Haralick Algorithmus GPU beschleunigt PRESA er Er a RR A y Ze yA E On O O gt gt abh ngig von Bild 4 3 Abh ngigkeitsgraph fir die Berechnung des Haralick Texturen Bildmerkmals in Kreise und Zwischenergebni
122. itt behandelt die Entwicklung von Kompilierern um eine geeignete Sprache in ein Hardwaredesign bersetzen zu k nnen 32 BOW Ne 3 3 Kompiliererentwicklung 3 3 1 Ubersicht Kompilierer Baukasten Die ersten Entwickler von Kompilierern mussten Assembler benutzen was viele Jahre Ent wicklungszeit ben tigte Von da an konnten weitere Kompilierer mit hohen Programmierspra chen schneller entwickelt werden Auch das war noch immer m hselig da die Komplexit t der Grammatiken beispielsweise in C sehr hoch ist Es wurden Programme entwickelt wie z B LEX die die lexikalische Analyse bernahm und YACC 45 die eine programmierte Gram matik syntaktisch analysiert und einen AST konstruiert Die Open Source Varianten hei en FLEX und BISON Der n chste Schritt der Entwicklung war es ein Framework zu program mieren das hilft Kompilierer zu entwickeln Es enth lt viele vorgefertigte modulare Struktur elemente eines Kompilierers zusammengefasst in eine Bibliothek mit Hilfsprogrammen zur Analyse Omniware 19 bietet ein Framework mit dem Fokus auf die sichere Ausf hrung von Programmmodulen in mobilen Endger ten an ohne die laufende Ausf hrungsumgebung sch digen zu k nnen Das Kompilierer Framework Phoenix 60 ist f r die Entwicklung von Kompilierer Backends gedacht die auf Microsoft Kompilierer aufsetzen ROSE 69 ist ein freies Kompilierer Framework das speziell f r objektorientierte Sprache geeignet ist LLVM ist derzeit
123. iziert Existiert zur Zeit der Aufnahme kein bin res Segmentbild aus einer anderen Quelle der Online Prozessierung kann ein Algorithmus der den Hintergrund entfernt fiir den FPGA tibersetzt werden Der vorgestellte Algorithmus 6 2 verwendet einen Schwellwert der tiber dem Rau schen liegt um den Hintergrund auf Null zu setzen __kernel void removeBGwithSementImg int multicell ints Z int w int idx get_global_id 0 get_global_id 1 int index idy w idx int idy Z index multicell index gt 10 multicell index 0 Quelltext 6 2 Kernelfunktion die ein Multizellbild einem Schwellwert unterzieht Die letzte Methode die alle Pixel unterhalb eines Schwellwertes auf Null setzt manipuliert sowohl die Pixel die sich im Hintergrund befinden als auch diejenigen die sich innerhalb der Zellen befinden Werden zu viele Pixel innerhalb der Zellen ver ndert bleibt die M glich keit einen komplexeren Segementierungsalgorithmus in OpenCL zu implementieren Kann die komplexe Kernelfunktion nicht mit dem FPGA Kompilierer bersetzt werden dann kann sie f r GPUs bersetzt werden ohne dass diese ver ndert werden m ssen 117 7 Fazit und Ausblick 7 1 Ziele der Arbeit Diese Arbeit bietet L sungen zur Online und Offline Prozessierung aus dem ViroQuant Pro jekt an Bei der Offline Prozessierung k nnen GPUs eingesetzt werden um einen Rechencluster zu beschleunigen Die Verarbeitung der Mikroskopb
124. jeder Takt einem Ergebnis eines Threads entspricht ben tigt die Kernelfunktion in Form einer Pi peline entsprechend viele Takte plus der Latenzzeit um alle Threads einer work group auszuf hren 5 4 8 Taktnetz Wie bereits erw hnt gibt es drei Taktnetze mit unterschiedlichen Frequenzen clkDDR3 200 MHz Benutzerschnittstelle Speicher clkUser 133 MHz Rahmendesign und Pipeline cIkPCle 125 MHz Benutzerschnittstelle PCIe Die Taktrate der Pipeline und weite Teile des Rahmendesigns miissen schneller getaktet sein als clkPCle weil es sonst zum Uberlauf der Datenpuffer in der PCIe Einheit kommen k nnte Ebenso sollte sie wegen eines m glichen Pufferiiberlaufs in der Speicherverwaltungseinheit nicht schneller als 200 MHz sein Die Komponenten der Pipeline sind fiir eine Frequenz mit 200MHz entwickelt worden Beliebige Kombinationen der Pipeline Komponenten hoher Ver drahtungsaufwand der Stall Logik der Datenflusskontrolle und der Speicherbus Oder Logik k nnen die Frequenz einschr nken Die Pipelinefrequenz wird auf vorsichtige 133 MHz ge 99 5 OpenCL zu FPGA bersetzer setzt Mit dieser Frequenz konnten eine Vielzahl unterschiedlicher Kernelfunktionen in synthe sef hige und lauff hige Pipelines bersetzt werden und garantieren so eine fehlerfreie berset zung im Zeitverhalten 5 5 OpenCL Laufzeitumgebung 5 5 1 FPGA Kommunikation 5 5 1 1 Entwicklungssteckkarte ML605 7 5 Bild 5 23 ML605 Entwicklungsp
125. ktion bergeben wird diese von der Funktion cicreateProgramWith Source eingelesen und einer Programm Struktur zugeordnet Die bersetzung passiert erst mit dem Aufruf von c1BuildProgran Befindet sich in der Programm Struktur die bersetzte Repr sentation wird diese in den FPGA geladen Handelt es sich um Quelltext wird dieser mit der Ausf hrung 1 eines Bash Skripts 2 des VHDL Kompilierers und 3 eines TCL Skripts 106 1 2 3 4 5 6 7 8 9 10 11 12 13 5 5 OpenCL Laufzeitumgebung in die bersetzte Repr sentation gebracht und dann in den FPGA geladen Bash Skript Es beinhaltet zwei Programmaufrufe und mehrere Pr fungen Als Parameter er wartet das Skript eine OpenCL Datei mit einer Kernelfunktion und der Dateierweitung cl Das Skript pr ft die Existenz des Parameters und die Existenz der Datei mit der cl Endung Das C Frontend wird aufgerufen clang c m32 emit llvm Si o Sfile bc um in die SSA Zwischensprache LLVM IR zu bersetzen Die Ausgabedatei hat eine bc Endung Der Optimierer reduziert die IR Befehle in ein kompaktes Programm opt 03 s o Sfile 11 file bc VHDL Kompilierer Dieser wurde im Abschnitt 5 3 ausf hrlich beschrieben Er bersetzt eine 11 Datei in eine pipeline vhdl Datei ssacompiler build pipeline 11 build pipeline vhd TCL Skript Die pipeline vhdl wird vom einem TCL Xilinx Skript Tool Command Lan guage in eine bit Datei ein Hardwarede
126. l 2009 Justin Richardson Kunal Gosrani Siddarth Suresh Jason Williams Alan D George Com putational Density of Fixed and Reconfigurable Multi Core Devices for Application Ac celeration RSSI 2008 Randi Rost John Kessenich Dave Baldwin The OpenGL Shading Language URL http www opengl org documentation glsl1 2002 Stephen C Johnson Yacc Yet Another Compiler Compiler Technischer Report 1979 Bernd Schwarz Jiirgen Reichardt VHDL Synthese Entwurf digitaler Schaltungen und Literatur 47 48 49 50 51 52 53 54 55 56 57 Systeme Oldenburg Verlag M nchen 2009 Andreas Kugel Guillermo Marcus und Wenxue Gao The MPRACE Framework URL http 1i5 ziti uni heidelberg de mprace Chris Lattner und Vikram Adve LLVM A Compilation Framework for Lifelong Program Analysis amp Transformation In CGO 04 Proceedings of the international symposium on Code generation and optimization Seite 75 IEEE Computer Society Washington DC USA 2004 doi 10 1109 CGO 2004 1281665 Chris Lattner und Vikram Adve LLVM Language Reference Manual 2011 G Lienhart G Marcus Martinez A Kugel und R Manner Rapid Design of Special Purpose Pipeline Processors with FPGAs and its Application to Computational Fluid Dynamics In Field Programmable Custom Computing Machines 2006 FCCM 06 14th Annual IEEE Symposium on Seite 301 302 april 2006 doi 10 1109 FCCM 2006 6
127. latine als PCle Steckkarte verwendet Das Bild 5 23 zeigt die Entwicklungssteckkarte von Xilinx 82 die in dieser Arbeit fiir den Kompilierer verwendet wird Die Karte ist mit einer Vielzahl von Bausteinen und Schnittstel len ausgestattet um m glichst viele unterschiedliche Beispieldesigns und Entwicklungsanwen dungen demonstrieren zu k nnen F r die Kompiliererentwicklung unwichtige Merkmale sind Gigabit Transceiver Gigabit Ethernet Schnittstelle USB 2 0 Schnittstelle ein DVI Monitor anschluss und ein 16 mal 2 Zeichen LCD Monitor Die wichtigen Hauptmerkmale in Bezug auf den Kompilierer die teilweise bereits erw hlt wurden sind hier aufgelistet 100 5 5 OpenCL Laufzeitumgebung DDR3 SO DIMM Speicher 512 MB Das 512 MB gro e Speichermodul wird als globaler Speicher verwendet Es ist m glich es durch ein gr eres zu ersetzen womit der Speicherkon troller neu konfiguriert werden muss PCI Express x8 Stecker Die PCIe Schnittstelle wird f r die Kommunikation zwischen Karte und Host PC verwendet Die Anbindung der DMA Engine liegt bei vier Lanes der ersten PCIe Generation PCIe x4 v1 0 200 MHz Differentialoszillator Dieser Oszillator ist die Taktquelle der Taktnetze clkUser und clkDDR3 Das Taktnetz clkPCle wird von der PCIe Schnittstelle gespeist LEDs und DIP Schalter Die Leuchtdioden werden f r die Fehleranzeige und der Betriebsan zeige verwendet Mit den Schaltern kann Testlogik aktiviert werden um die o
128. lge auf siehe Bild 5 8 SourceAnalyser analyse analyseParallelStruct analyseParallellndices analyseBranchPart SourceParser analysePipelinePart analyseTransferPart analyseMemAccessPattern Bild 5 8 Auszug Klassendiagramm um die Analyseklasse analyseParallelIndices untersucht welche parallele Indizierung im Quelltext vorhanden ist Das Untersuchungsergebnis ist welche Dimensionierung die Parallelindizes benutzen Die Analysemethode sucht im Parsebaum nach ca11 Instruktion hinter der sich folgende Funkti onsaufrufe 35 befinden k nnen e uint get_work_dim Gibt die Anzahl benutzter Dimensionen zur ck size_t get_global_size uint D Dimensionsgr e der globalen work items e size_t get_global_id uint D Globale work item Nummer der entsprechenden Dimen sion size_t get_local_size uint D Dimensionsgr e der lokalen work items size_t get_local_id uint D Lokale work item Nummer der entsprechenden Dimensi on size_t get_num_groups uint D Dimensionsgr e der work groups size_t get_group_id uint D Nummer der work group der entsprechenden Dimension Mit der Abfrage des Parameters n kann die Dimension der entsprechenden Gruppe ausgelesen werden Die Dimensionierung gibt Hinweise auf das Zugriffsmuster der Speichertransfers z B das Lesen eines Arrays einer Matrix oder eines Volumens Die Methode analyseParallelstruct ist ohne Implementierung Sie ist ein Pl
129. lt die eine Art zweidimensionale Histogramme darstel len Auf den Matrizen werden die Bildmerkmale berechnet Die Matrizen und die Bildmerk male werden im Folgenden beschrieben und f r eine beschleunigte Berechnung analysiert 4 1 1 Co occurrence Matrizen Die Berechnung der Co occurrence Matrizen Co Matrizen basiert auf einer Statistik zweiter Ordnung und ist in 37 und 38 beschrieben Es werden Histogramm Matrizen Co Matrizen anhand benachbarter Pixelpaare erstes und zweites Pixel aus dem Quellbild gebildet Die Pi xelpaare werden f r eine bestimmte Sichtweise angeordnet gemeint ist ein bestimmter Pixel abstand und ein Winkel des ersten und des zweiten Pixel zueinander Wobei der Grauwert des ersten Pixel die Zeile der Co Matrizen adressiert und das zweite Pixel die Spalte Die ermittelte Zelle wird entsprechend eines Histogramms um Eins akkumuliert Die Co Matrix ist vollst n dig wenn alle Pixelpaare die aus dem Quellbild entstehen k nnen aufsummiert worden sind Ein ausf hrliches Beispiel gibt es in 37 F r jede Sichtweise wird eine eigene Co Matrix gebildet die ein Vorkommen der Grauwerte unter einer gewissen Anordnung der Pixelpaare repr sentiert Veranschaulicht sind die Matrizen eine Kombination aus Ortsinformationen und zweidimensionalem Histogramm Die Gr e der Co Matrizen ist abh ngig von der Anzahl an m glichen Grauwerten im Quell bild Bei einer Bittiefe von 12 Bit existieren 4096 unterschiedliche Grauw
130. m done Signal die Adresse die Daten und die Information ob es sich um einen Burst Zugriff handelt senden Das valid Signal zeigt ob der Bus f r einen Transfer beansprucht wird Auf dem Lese Bus gibt es die Problematik dass der Slave der Daten angefordert hat wieder identifiziert werden muss um die Daten richtig zuzuordnen Gel st wird das Pro blem mit einem Job Puffer und einer Zustandsmaschine die Speicheranfragen mit der Slave ID zwischenspeichert Die Anfrage besteht aus der Adresse einem Burst Signal 5 4 Rahmendesign clkUser rst ddr3_rd_valid ddr3_rd_data dma_wr_en dma_wr_addr dma_wr_data dma_rd_addr dma_rd_burst dma_rd_en bus_rd_burst bus_rd_addr_valid bus_rd_addr bus_wr_addr bus_wr_data bus_wr_valid bus_wr_burst bus_wr_write bus_ctrl_rd_done bus_ctrl_wr_done bus_def_rd_size bus_def_wr_size ctrl_pr_config ctrl_block_end ctrl_block_start MemManager FIFO Zustandsmaschine ddr3_rd_addr ddr3_rd_burst ddr3_rd_en ddr3_wr_addr ddr3_wr_burst ddr3_wr_en ddr3_wr_valid ddr3_wr_data dma_rd_data dma_rd_valid bus_rd_data bus_rd_data_valid bus_rd_data_sel bus_ctrl_rd_sel bus_ctrl_wr_sel Bild 5 21 Schematisches Blockschaltbild der Speicherverwaltungseinheit 97 5 OpenCL zu FPGA bersetzer und einem Giiltigkeitssignal ob es eine Leseanforderung gibt Kommen nach der La tenzzeit des Speicherkontrollers die Daten wird aus dem Job Puffer di
131. menge aller OpenCL Funktionen in einer Biblio thek implementiert genug um eine Demonstration zu erm glichen Die nicht implementierten Funktionen geben einen Hinweis aus und beenden das Programm falls sie von einer OpenCL Anwendung verwendet werden Die OpenCL Funktionen der Laufzeitumgebung sind ANSI C w hrend die Kommunikations klasse und die MPRACE Bibliothek in C entwickelt sind Das bedeutet dass aus den C 104 on Dn Fw NY eS 5 5 OpenCL Laufzeitumgebung Funktionen der Laufzeitumgebung Klassenmethoden aufgerufen werden miissen Abhilfe schaf fen R ckruffunktionen call back functions die in der C Umgebung deklariert aber in C definiert werden Die Implementierung der Call Back Funktionen erstellt die n tigen Klassen Instanzen und setzt die C Funktionsaufrufe in Klassen Methodenaufrufe um Es existiert f r je de Methode der Kommunikations Klasse DeviceTrans und der Speichertabellen Klasse Device Mem eine entsprechend bezeichnete Call Back Funktion Die Funktion systemca11 ist eine selbst entwickelte Variante der Bibliotheksfunktion system die Linux Programmaufrufe ausf hrt Die Bibliotheksfunktion lie sich nicht korrekt aus f hren weil deren fork zur Erzeugung eines Kind Prozesses eine Kopie der DMA Puffer erstellt der physikalische Speicheradressen enth lt Sobald der Kind Prozess beendet wird werden alle Speicherbereiche inklusive DMA Puffer freigegeben auch der physikalische Spei cher
132. mmable Logic and Applications 2005 International Conference on Seite 317 322 aug 2005 doi 10 1109 FPL 2005 1515741 75 Justin Tripp Preston Jackson und Brad Hutchings Sea Cucumber A Synthesizing Com piler for FPGAs In Manfred Glesner Peter Zipf und Michel Renovell Editors Field Programmable Logic and Applications Reconfigurable Computing Is Going Mainstream Volume 2438 von Lecture Notes in Computer Science Seite 51 72 Springer Berlin Hei delberg 2002 doi 10 1007 3 540 46117 5_90 76 Richard Wain lan Bush Martyn Guest Miles Deegan Igor Kozin Christine Kitchen Cheshire Wa Ad Richard Wain lan Bush Martyn Guest Miles Deegan Igor Kozin und Christine Kitchen An overview of FPGAs and FPGA programming Initial experiences at Daresbury Technischer Report 2006 doi DL TR 2006 010 77 Xiaojun Wang und Miriam Leeser VFloat A Variable Precision Fixed and Floating Point Library for Reconfigurable Hardware ACM Trans Reconfigurable Technol Syst Volume 3 16 1 16 34 September 2010 doi 10 1145 1839480 1839486 78 Markus Wannemacher Das FPGA Kochbuch MITP Verlag 1998 79 Xilinx Development System Reference Guide v10 1 Edition 2008 80 Xilinx Partial Reconfiguration User Guide ug702 v12 3 Edition Oktober 2010 81 Xilinx Virtex 6 FPGA Integrated Block for PCI Express User Guide 2010 82 Xilinx Virtex 6 FPGA ML605 Evaluation Kit URL http www xilinx com products
133. mmierer entwickelt worden die eine einfache C hnliche Sprache gegen ber den g ngigen Hardware Beschreibungssprachen bevorzugen Im Benutzerhandbuch 42 sind alle nicht standardisierten C Spracherweiterungen erl utert die eine gezielte Hardwaresynthese auf einer h heren Abstraktionsebene erm glichen Die Konzepte von Handel C basieren auf seriellen und parallelen Programmabl ufen und Kanal Kommunikation Das Bild 3 2 demonstriert die Konzepte und zeigt Details der Sprache void main void Block 1 int 2 a int 4 b int 8 c int 8d int 10 e wann chan ch with fifolength 8 b afa Block 1 Block par 2 seq y b b ta Block 2 c b 3 Block 3 ch e Kanal 1 3 Block _ Kanal 1 Block 4 ch d d dq gt gt 2 Block 4 Block e c d Block 5 5 a Handel C Quelltext b Strukturdiagramm Bild 3 2 Handel C Quelltextbeispiel a mit Strukturgraph b so angeordnet dass die parallelen Abl ufe deutlich werden e In Handel C werden Integer Variablen mit einer frei w hlbaren Bitbreite angegeben 36 3 4 Software Hardware Kompilierer Weiter hat man die Wahl zwischen Vorzeichen behaftetem und Vorzeichen losem In teger e Die Ausf hrung wird in serielle Abl ufe seq in b von oben nach unten und in parallele Istinlinelparl in b auf gleicher Ebene unterteilt Im parallelen Segment gibt es zwei serielle Abl ufe e Jede Zuweisung entspricht einer Taktst
134. mpilierer 5 3 1 bersetzungskette Es gibt viele Ans tze wie der OpenCL Quelltext bersetzt werden kann Beispielsweise l sst sich mit Lex und Yacc ein OpenCL Frontend entwickeln das den Quelltext in einen Parser baum im Speicher umsetzt Diese Aufgabe ben tigt viel Entwicklungszeit Statt dessen k nnte man das C Frontend von GCC verwenden dann br uchte man keinen C Parser entwickeln und man k nnte die Entwicklung auf die Zwischensprache aufsetzen Da das GCC C Frontend in die Jahre gekommen ist wurde es mit LLVM berarbeitet und erneuert Auch diese Arbeit setzt auf LLVM auf Das LLVM C Frontend Clang bersetzt in eine standardisierte Zwischenspra che f r die es in dieser Arbeit einen Parser und ein Backend f r den FPGA zu entwickeln gilt Die bersetzungskette l sst sich in drei Glieder aufteilen die in Bild 5 2 angeordnet sind Das Clang Frontend wird f r die bersetzung der Kernelfunktion in LLVM IR verwendet Besser w re es ein OpenCL Frontend zu verwenden da die Kernelfunktion OpenCL spezifi sche Schl sselworte enthalten kann die von Clang nicht identifiziert werden Beispielsweise spezifizieren die Schliisselworte 1oca1 und globai die Speicherbereiche in denen die Variablen 67 5 OpenCL zu FPGA bersetzer OpenCL Kernel bersetzung Clang Frontend llvm opt Vhdl Backend Bild 5 2 Ubersetzungskette des VHDL Kompilierers abgelegt werden Folglich kann das Clang Frontend ke
135. n Transferlesebl cken und den Transferschreibbl cken existieren um sie selektieren zu k nnen 89 5 OpenCL zu FPGA bersetzer Pipeline Kennung Es ist ebenso wichtig zu wissen welche Pipeline sich gerade in der Hard ware befindet Deswegen muss jedes Design ber eine ID verf gen Eine Zufallszahl zwischen 0 und 255 ist ausreichend zuf llig genug um eine generierte Pipeline zu kennzeichnen 5 3 7 Einschr nkungen in der bersetzung Der bersetzungsvorgang von einer OpenCL Kernelfunktion in eine VHDL Pipeline unter liegt Einschr nkungen die hier aufgelistet sind Der Shared Speicher kann nicht verwendet werden deshalb ist eine Threadkommunika tion nicht m glich Schleifen werden nicht unterst tzt Es fehlt die Implementierung der Sprungbefehle und deren bersetzungsvorschrift die im Prototyp nicht vorgesehen sind Komplexe is Strukturen k nnen ebenfalls wegen fehlender Sprungbefehle nicht ber setzt werden Auch dieser Punkt ist Ziel der n chsten Version Datenabh ngige Speicherzugriffe und komplexere Speicherzugriffsmuster werden nicht unterst tzt Die Datenbreite der Pipelinearchitektur ist f r 32 Bit entwickelt worden F r andere Bitbreiten muss die Entwicklung angepasst werden Bisher gibt es keine Implementierung der Flie kommaoperatoren in den Rechenbl cken F r zwei der Einschr nkungen gibt es Behelfsl sungen die weniger effizient sind Kleinere Schleifen k nnen mit
136. n Das stall Signal zeigt dass der Datenpuffer leer ist und gibt der Pipeline den Be fehl die Berechnungen anzuhalten Der in Bild 5 16 dargestellte Transferschreibblock hat die Aufgabe Ergebnisdaten aufzuneh men in einem Datenpuffer zwischenzuspeichern und sie ber den Bus in den Speicher zu transferieren BLOCK_NUM BlockTransferWr clkUser gt acalc_next rst gt acalc_update acalc_addr acalc_request acalc_rdy bus_ctrl_wr_done acalc_burst Zustandsmaschine bus_wr_addr acalc_valid bus_wr_burst acalc_fin bus_wr_write bus_ctrl_wr_sel bus_wr_data ee bus_wr_valid data_valid data_stall Bild 5 16 Schematische Darstellung des VHDL Transferschreibblocks Oben Generic Ports und unten die Entity Ports 86 5 3 VHDL Kompilierer e Die Ergebnisse kommen ber die data Schnittstelle im 32 Bit Format an und werden in den Datenpuffer geschrieben Wenn er voll ist gibt das stall Signal der Pipeline die Anweisung anzuhalten e Vergleichbar mit der Zustandsmaschine im Transferleseblock hat diese ebenfalls die Kontrolle ber die acalc Schnittstelle Wenn der Datenpuffer Daten enth lt wird die Adresse angefordert an der die Daten gespeichert werden sollen e ber die bus Schnittstelle gibt die Zustandsmaschine den Schreibbefehl zusammen mit den Daten auf den Speicherbus Der Datenpuffer ordnet die 32 Bit Ergebnisse in 256 Bit breite Datenworte um markiert die g ltigen 32 Bit Werte und f llt wenn n tig fehlende
137. n Die Ausf hrung einzelner Operationen unterliegt allerdings Regeln denn eine Instrukti on wird parallel auf mehreren TPs nach dem SIMD Prinzip ausgef hrt Warp Scheduler Ein Scheduler ist eine Ausf hrungseinheit die Instruktionen aus dem Spei cher lie t dekodiert und auf die TPs verteilt Diese Verteilung passiert in Hardware was bli cherweise ein Betriebssystem in Software erledigt Ein Warp ist eine Gruppe von 32 Threads Ausf hrungsf den D h dass ein Warp Scheduler eine Instruktion auf 32 TPs parallel mit unterschiedlichen Daten ausf hrt In der Hardware werden allerdings lediglich 16 TPs das entspricht einem half Warp beauftragt mit der gleichen Instruktion zu rechnen Erst zu einem sp teren Zeitpunkt wird die zweite H lfte des half Warps ausgef hrt Damit zu jedem Zeitpunkt 2 Grundlagen alle 32 TPs einer SM besch ftigt sind gibt es zwei Warp Scheduler die zu einem Zeitpunkt zwei half Warps unterschiedlicher Warps ausf hren Lade und Speichereinheiten IO Jede SM besitzt 16 Lade und Speichereinheiten die pro Takt ein Ziel oder eine Quelle im Speicher adressieren k nnen und einen Speichertransfer in itiieren Der Speichertransfer richtet sich als erstes an den L1 Cache bevor die Daten vom oder zum Massenspeicher transferiert werden Spezielle Funktionseinheiten SFU In ihnen werden komplexere Instruktionen abgearbeitet die die TPs nicht unterst tzen Beispiel f r die Instruktionen sind s
138. n Some Computer Organizations and Their Effectiveness IEEE Trans Comput Volume C 21 948 1972 doi 10 1109 TC 1972 5009071 W Gao und Z Han PCIe SG DMA controller URL http opencores org project pcie_sg_dma 2011 Wenxue Gao Andreas Kugel Reinhard M nner und Guillermo Marcus PCI Express DMA Engine Design Technischer Report CBM Progress Report 2007 M Gipp G Marcus N Harder A Suratanee K Rohr R K nig und R M nner Ha ralick s Texture Features Computations Accelerated by GPUs in Biological Applications In GPU Technology Conference GTC 09 NVIDIA San Joes California USA 2009 Poster Research Summit Markus Gipp Guillermo Marcus Nathalie Harder Apichat Suratanee Karl Rohr Rainer K nig und Reinhard M nner Haralick s Texture Features using Graphics Processing Units GPUs In Proceedings of The World Congress on Engineering 2008 Volume I von ICPDC 08 Seite 587 592 International Association of Engineers Newswood Limited London UK UK 2008 Markus Gipp Guillermo Marcus Nathalie Harder Apichat Suratanee Karl Rohr Rai ner K nig und Reinhard M nner Haralick s Texture Features Computed by GPUs for Biological Applications LAENG International Journal of Computer Science Volume 36 129 Literatur 34 ja 35 a 36 37 38 39 40 41 42 a 43 44 45 46 130 2009 Markus Gipp Guillermo Marcus Nathal
139. n ausl sen Die Ausgabesignale m ssen w hrend des Programmierens von der statischen Logik entkoppelt werden Der einfachste Weg der Ent koppelung ist alle Signale mit der Quelle aus dem dynamischen Modul im statischen Design mit einem UND Gatter zu verkn pfen Mit einem Steuersignal kann der Ausgang des UND Gatters und somit das Signal aus dem dynamischen Design gezielt auf logisch Null gesetzt werden Das UND Gatter bezeichnen wir als Torschaltung das von einem Kontrollsignal ge steuert wird und das Tor ffnen und schlie en kann Wie die OpenCL Funktion c1BuildProgram die Kernelfunktion bersetzt wurde bereits be schrieben Es fehlt wie das generierte dynamsiche Pipelinemodul in den FPGA geladen wird 1 Die Bit Datei mit dem Pipelinemodul wird vom Dateisystem gelesen die Dateil nge bestimmt und in eine OpenCL Programm Struktur gespeichert 2 Mit der Call Back Funktion devsetconfigModerr 1 wird ein Bit im Kontrollregister ge setzt und die Torschaltung schlie t Die Signale aus dem dynamischen Modul sind im statischen Design logisch entkoppelt 3 Das Pipelinemodul wird mit devConfigurePR program gt bitfile program gt bitlength ber die ICAP Schnittstelle in den FPGA geladen 4 Danach wird mit devsetconfigModepr 0 die Torschaltung ge ffnet und die Signale aus der Pipeline haben wieder Einfluss auf die statische Logik 5 Ein Pipelinereset mit der Call Back Funktion devresetPipeline beendet den Austausc
140. n m ssen Die dritte Variante ist die Oderlogik bei der alle Quellsignale mit Odergatter zu einem Signal kom biniert werden Dabei m ssen alle Quellsignale im inaktiven Zustand ein Low Signal ausgeben und nur der aktive Teilnehmer darf beide Signalzust nde verwenden Die Oderlogik hat an die ser Stelle das beste Zeitverhalten und wurde f r diese Arbeit gew hlt Stall Logik und Datenfluss Kontrolle Die Oderlogik kommt auch zum Einsatz um die Stall Signale zu einem zu kombinieren Sofern kein VHDL Block einen Datenmangel Transferle seblock bzw einen Datenstau Transferschreibblock aufweist kann das kombinierte Stall Signal verwendet werden um den Transferlesebl cken im n chsten Takt der Pipeline einen Wert zu liefern Diese Logik steuert den Datenfluss der Pipeline Registerstufen im Speicherbus Das Zeitverhalten erfordert Registerstufen zwischen dem Rahmendesign und dem Pipelinedesign um die Speicherbussignale in beide Richtungen zeit lich zu entkoppeln Diese Ma nahme erm glicht den Einsatz einer dynamischen partiellen Rekonfiguration um Pipelinemodule austauschen zu k nnen Pipeline Fertigsignal Die Pipeline muss zeigen wann sie mit allen Berechnungen fertig ist Dieses Signal wird dem Adressierungsblock entnommen denn dieser kennt die Adresse vom letzten Datenelement und wei wann sie f r einen Datentransfer verwendet wurde Anzahl der Transferinstanzen Das Rahmendesign muss wissen wie viele Instanzen von de
141. n und bestimmt 27 3 Stand der Technik welche Co Prozessoren f r diese Arbeit in Frage kommen CPU Die CPU Central Processing Unit hat einen umfassenden Befehlssatz mit dem sich Programme entwickeln lassen die keine Einschr nkung in ihrer Funktionalit t haben Eine Cache Hierarchie reduziert die Zugriffszeiten auf den Hauptspeicher die viele hundert Pro zessortaktzyklen dauern k nnen Die Taktraten von CPUs um die 3GHz steigen kaum noch Das mooresche Gesetz eine Verdoppelung der Integrationsdichte bei integrierten Schaltkreisen ICs alle 18 bis 24 Monate wird genutzt um die Architektur einer CPU in ihrer Leistungs f higkeit zu steigern Alle modernen CPUs besitzen eine mehrstufige Pipeline um in prak tisch jedem Takt eine Instruktion auszuf hren Eine Weiterentwicklung geht sogar dahin dass mehrere Instruktionen in einem Takt ausgef hrt werden k nnen Ein Beispiel stellt die Hy perThreading Technik von Intel da die einen virtuellen Kern bietet um in einem Takt mehrere Instruktionen ausf hren zu K nnen Eine andere Technik den Grad der parallelen Ausf hrung zu erh hen bieten die Streaming SIMD Extensions Register SSE Mit ihnen k nnen einzel ne Instruktionen auf mehrere Daten angewendet werden Seit vielen Jahren steigt die Anzahl der SSE Register deren Registerbreite und der SSE Instruktionen an um immer mehr Daten nach dem SIMD Prinzip parallel zu verarbeiten Die letzte Steigerung der Integrationsdichte li
142. nbus anzuschlie en daf r keine Ca ches in den Transfer Bl cken zu installieren hnlich wie bei einer CPU Gibt es viele Transfer bl cke in der Pipeline wird der Datenbus zum Flaschenhals da jeder Transferblock gleich zeitig Daten aus dem Cache anfordert Aus diesem Grund werden viele Caches eingesetzt um jedem Transferblock Zeit zu geben Daten aus dem Speicher anzufordern w hrend die Caches die Pipeline mit Daten versorgen Der Busmaster der den Datenbus kontrolliert wird in der Speicherverwaltungseinheit des Rah mendesigns implementiert 5 4 5 Speicherkontroller und vereinfachtes Ansprechen Um Zugriff auf das 512 MByte gro e DDR3 Speichermodul zu bekommen das auf der Ent wicklungssteckkarte gesteckt ist ben tigt der FPGA einen Speicherkontroller Der Virtex6 be sitzt keinen Hardcore Speicherkontroller somit muss ein Softcore beschrieben werden Diese Arbeit nutzt das Xilinx MIG 83 Memory Interface Generator um eine Hardwarebeschrei bung eines Speicherkontrollers 85 erzeugen zu lassen Die Benutzerschnittstelle des erzeugten Speicherkontrollers l sst sich logisch in drei Teile glie 94 5 4 Rahmendesign dern Kommando Schnittstelle Mit ihr kann man zusammen mit einer Adresse Lese und Schreib Kommandos absetzen Gleichzeitig gibt es Signale die mitteilen dass der Spei cherkontroller gerade keine Kommandos akzeptiert beispielsweise wenn gerade ein Re fresh Zyklus passiert Sc
143. ne Datenbank Native Generic Database die den Logikschaltplan auf einfache Logikelemente wie UND Gatter ODER Gatter Deko dierer Flip Flops und RAMs herunter bricht und auflistet Neben dem Logikschaltplan enth lt die NGD Datei auch die Logikbeschreibung 3 MAP bildet die einfachen Logikelemente auf die Logikzellen eines speziellen FPGAs ab so dass die CLBs m glichst effizient genutzt werden Das Werkzeug liefert eine NCD Datei Native Circuit Description 4 PAR Place And Route platziert und verdrahtet die vorkonfigurierten CLBs im FPGA Wieder liefert das Werkzeug eine NCD Datei die diesmal die ben tigten Chip Ressour cen nicht nur listet sondern auch mit Ortsinformationen und Verbindungswegen versieht 5 BitGen bringt die platzierten und verdrahteten CLBs in ein Bin rformat das f hig ist den FPGA ber die Konfigurationsschnittstelle zu programmieren Das Ergebnis der bersetzung ist eine BIT Datei Bitstream file mit der urspr nglichen Hardwarebe schreibung Wird auch nur ein Logikgatter in der Hardwarebeschreibung ver ndert muss der gesamte ber setzungsprozess neu angesto en werden Dieser Nachteil l sst sich mit einer Partitionierung 20 2 3 Kompiliererentwicklung eingrenzen Die bersetzungsschritte 1 4 k nnen f r jede Partition unabh ngig geschehen Die nderung einer Partition mit dem Durchlauf der bersetzungskette nutzt die bereits bersetz ten anderen Partitionen f r die Platzier
144. ner Zwischensprache Der AST wird in diesem Schritt in eine Zwischenspra che bersetzt die der Zielsprache hnelt Optimierung der Zwischensprache Wegen der standardisierten Zwischensprache gibt es vie le Optimierer die den Quelltext im Umfang reduzieren Z B Zwischenergebnisse werden in Registern gehalten ohne sie neu zu berechnen oder unerreichbarer Quelltext wird entfernt Generierung der Zielsprache Erst jetzt wird aus der Zwischensprache die Zielsprache gene riert Diesen Teil nennt man Assembler Dank des bereits optimalen Quelltextes in der Zwi schensprache l sst sich der Assembler f r andere Computerarchitekturen austauschen Der Assembler modifiziert den Quelltext f r dessen Architektur Beispielsweise wird die Regis teranzahl und die Registerbreite in der Zielsprache angepasst Das Frontend und das Backend sind sehr ausf hrlich in Compiler Design in C 41 beschrie ben 22 3 Stand der Technik 3 1 Co Prozessoren 3 1 1 Beschleunigung Co Prozessoren werden eingesetzt um die CPU zu entlasten und somit die gesamte Ausf h rungszeit eines Programms zu reduzieren Dabei wird der rechenintensive Teil bzw der Teil der besonders viel Ausf hrungszeit beansprucht vom parallel arbeitenden Co Prozessor ber nommen Die gesamte Ausf hrungszeit eines Programms l sst sich in zwei Teile zerlegen einen der nicht beschleunigt wird ts und einen der beschleunigt bzw parallelisiert wird tp Welche Be
145. ng ben tigt Mit dem Index der gepackten Co Matrix wird der zugeh rige Grauwert ermittelt Die Grauwert Indextabelle wird bei der Co Matrixgenerierung ben tigt F r den Grauwert der Zel le wird der Index der gepackten Co Matrix nachgeschlagen Die Grauwerte Index Zuordnung f r eine Zelle ist eine 1 1 Beziehung Ermittelt wird die Zu ordnung indem alle vorkommenden Grauwerte in einem Vektor markiert werden und die Zei len der unmarkierten Grauwerte entfernt werden Die Gr e des brig gebliebenen Vektors entspricht der Anzahl der vorkommenden Grauwerte aus der Zelle Aus der eben gewonnenen Index Grauwerttabelle wird durch Umkehrung die Grauwert Indextabelle gewonnen 59 4 Haralick Algorithmus GPU beschleunigt 4 3 2 3 Gepackte Co Matrix gezielt generieren Wie viele Grauwerte in einer Zelle tats chlich vorkommen ist von der Generierung der Lookup Tabellen bekannt Die Anzahl der vorkommenden Grauwerte ist gleich der Kantenl nge der zu generierenden gepackten Co Matrix Somit wird die gepackte Co Matrix gezielt generiert oh ne eine volle Co Matrix im Speicher halten zu m ssen Nachdem ein Grauwertepaar gebildet wurde werden dessen Grauwerte mit der Grauwert Indextabelle durch Indices der gepackten Co Matrix ersetzt Das identifizierte Element der gepackten Co Matrix wird um eins erh ht hnlich wie bei einem Histogramm Bei der Adressierung einer Datenreihe im Grafikkartenspeicher gibt es f r effiziente Speicher zugriff
146. nos org conformance adopters conformant products 2011 Hardware description language URL http de wikipedia org wiki HDLC 2011 Learn About Altera s OpenCL Program for FPGAs URL http www altera com b opencl html 2011 LLVM Users URL http 11vm org Users html 2011 Intel microprocessor export compliance metrics URL http www intel com support processors xeon sb CS 020863 htm 5 Dez 2008 Ali Reza Adl Tabatabai Geoff Langdale Steven Lucco und Robert Wahbe Efficient and language independent mobile programs In Proceedings of the ACM SIGPLAN 1996 conference on Programming language design and implementation PLDI 96 Seite 127 136 ACM New York NY USA 1996 doi 10 1145 231379 231402 Gene M Amdahl Validity of the single processor approach to achieving large scale computing capabilities In Proceedings of the April 18 20 1967 spring joint computer conference AFIPS 67 Spring Seite 483 485 ACM New York NY USA 1967 doi 10 1145 1465482 1465560 P Banerjee D Bagchi M Haldar A Nayak V Kim und R Uribe Automatic conver sion of floating point MATLAB programs into fixed point FPGA based hardware design In Field Programmable Custom Computing Machines 2003 FCCM 2003 11th Annual IEEE Symposium on Seite 263 264 april 2003 doi 10 1109 FPGA 2003 1227262 Michael J Beauchamp Scott Hauck Keith D Underwood und K Scott Hemmert Em bedded floating point uni
147. nsertSignalDeclarationBlock compile insertSignalDeclarationConst insertComponentinstances setBlockSizeRd setBlockSizeWr Bild 5 13 Ausschnitt Klassendiagramm betreffend der VHDL Generierung Im Gegensatz zu den vorherigen Klassenbeziehungen ist die Bindung dieser Klasse an die SourceCompiler Klasse gering Die sourceCompiler compile Methode erzeugt eine vnalarchi tecture Instanz die ausschlie lich mit string Datentypen Textbausteine der Klasse hinzuf gt hnlich wie die tostring Methoden in Javaklassen haben auch die VHDL Block Instan zen die Methode getvhdlInstance die eine VHDL Textrepr sentation der Komponente zu r ck gibt Der VHDL Text jedes VHDL Blocks im Block AST wird ber die insertcomponent Instances hinzugef gt Mit der Methode insertsignalDeclarationBlock werden der VHDL Klasse alle VHDL signale mitgeteilt die von den VHDL Bl cken gebraucht werden Es gibt mit der Methode insertsignalDeclarationConst die Option weitere Signale hinzuzuf gen 83 5 OpenCL zu FPGA bersetzer die von der zus tzlichen Logik siehe folgenden Abschnitt verwendet werden Die gesam melten Textbausteine der Signaldeklarationen der Blockinstanzen und der zus tzlichen Logik werden in der Methode writevnalcode in eine VHDL Datei geschrieben 5 3 6 2 VHDL Bl cke als Bausteine Bild 5 14 zeigt den funktionalen Aufbau des Adressierungsblocks Er hat die Aufgabe die Adressen zu berechnen d
148. nteger Arithmetik Logik Aufgab en und Bit Operationen Die ideale Anwendung m sste eben aus einer Pipeline bestehen mit genau den Operationen die der FPGA gut kann Der FPGA erzielt besonders hohe Rechen geschwindigkeit dadurch dass in jedem Takt alle Operationen der Pipeline parallel ausgef hrt werden Wenn die Abw rme bzw Verlustleistung ein Problem darstellen werden FPGAs be vorzugt eingesetzt Eine Untersuchung 43 von Williams zeigt dass FPGAs einen besseren Wirkungsgrad gemeint ist Rechenleistung pro Watt als GPUs haben Der FPGA ist bei Inte geroperationen um das sechsfache und bei Flie kommaoperationen um das zweifache effizien ter im Energieverbrauch f r die gleiche Rechenleistung Ein FPGA ben tigt ca 25W Leistung hat aber die Nachsicht bei den Anschaffungskosten von 1000 10000EUR Die Kosten ver einen die Entwicklung oder die Anschaffung einer Steckkarte mit einem FPGA Chip Eine Aussage ber die GFLOPS zu treffen ist nicht einfach Ein Virtex 5 mit 330 Tausend Logik zellen konnte laut 76 56 GFLOPS erreichen einschlie lich aller arithmetischen Operatoren aus den DSP slices und den vorhandenen Logigbl cken Inzwischen gibt es Virtex 6 FPGAs und Virtex 7 sind angek ndigt diese laufen mit h heren Taktraten und haben mehr Logikzel len F r die neuste Generation wird die Rechenleistung auf 100 GFLOPS gesch tzt Es ist uns nicht wichtig eine genaue Zahl zu berechnen da es sich lediglich um die theoretisch maxi
149. onen k nnen den kleinen Speicherbereich lediglich lesen Dieser Speicherbereich hat an Be deutung verloren seit dem der globale Speicher ber ein Cachesystem verf gt Denn dieser besitzt einen separaten Cache um m glichst immer gute Cachetreffer zu erzielen Texturenspeicher Dieser ist vergleichbar mit dem konstanten Speicher der ber ein separates Cache verf gt Allerdings kann er wesentlich gr er sein indem Teile des globalen Speichers f r den Texturenzugriff reserviert werden Wenn die Kernelfunktion viele Ressourcen ben tigen kann die Effizienz leiden Braucht ein CUDA Block den gesamten Shared Speicher kann somit nur dieser auf einer SMs ausgef hrt werden statt m glichen acht Die Effizienz leidet auch wenn wenige Threads zu viele Register ben tigen dann k nnen eben auch nur wenige Threads gleichzeitig auf einer SM ausgef hrt werden Die limitierenden Faktoren sind in Formeln einer Excel Tabelle hinterlegt die die Aus 12 So oa wu WwW WN e U N 2 1 Grafikkarten als Rechenbeschleuniger lastung der SMs bzw die Anzahl parallel arbeitender CUDA Bl cke berechnet siehe 66 Die typische Abfolge eines einfachen CUDA Programms beginnt mit der Kopie der Eingangs daten in den Grafikkartenspeicher Beim Kernelaufruf wird die Anzahl an Bl cken und Threads definiert und dimensioniert je nachdem wie h ufig es gew nscht ist das Kernelprogramm pa rallel auszuf hren Nach der Berechnung werden die Ergebni
150. onsdichte der zuk nftigen Chips steigert bei FPGAs die Anzahl und die Komplexit t der Logikzellen um einen h heren Grad der Parallelisierung zu erreichen Die Frequenzen der FPGAs steigen nur zweitrangig an OpenCL ist heute und in der Zukunft eine ideale Sprache Algorithmen wahlweise parallel auf Vielkern CPUs GPUs oder FPGAs ausf hren zu lassen Die Ergebnisse dieser Arbeit konnten teilweise in Journals einem Konferenzbeitrag und in einem Poster ver ffentlicht werden Die Ver ffentlichung zur Beschleunigung der Haralick Bildmerkmale auf der GPU wurde auf der ICPDC 08 zum Best Paper normiert 32 Auf ei ne Einladung hin wurde der Konferenzbeitrag zu einem Journal erweitert 33 und im JAENG International Journal of Computer Science ver ffentlicht Die Ergebnisse der zweiten GPU Version konnten im Proceedings of the Fourth International Conference on High Performan ce Scientific Computing vorgestellt werden 34 Zuletzt wurden technische Details zur GPU Beschleunigung auf der NVIDIA Konferenz GTC 09 mit einem Poster dem Fachpublikum pr sentiert 31 7 2 Verbesserungen f r die Zukunft Wie bereits erw hnt lohnt sich eine weitere Beschleunigung der Haralick Bildmerkmalserken nung nicht und es besteht auch kein Bedarf einer Weiterentwicklung Der OpenCL FPGA Kompilierer besitzt den Stand eines Prototyps der eine gezielte Aufga be erf llt F r eine allgemeinerer Verwendung wurde explizit auf eine modulare Entwicklun
151. penRCL Implementierung ihre Geschwindigkeitsleistung erzielt Die Analyse ist neben dem Kernel Scheduler die entscheidende nderung zur klassischen OpenCL bersetzungs kette Die Generierung der Zielsprache und das Linken wurde auf ein existierendes LLVM Backend aufgesetzt und der speziellen MIPS Multithread Archtitektur auf dem FPGAs ange 42 3 4 Software Hardware Kompilierer E externe i A Bibliothek E FORTRAN i cpu LV ee gt M Generierung I bersetzte und linken der Kernel OpenCL Objekte Zielsprache Funktionen Ea Speicher lt A PESA a Kernel gcc Kernel 2 zugriff arth Scheduler Front end Optimierung Bild 3 7 Leicht modifizierte OpenCL Ubersetzungskette Die grauen Bl cke machen den Un terschied von der OpenRCL Implementierung aus passt Die Main Funktion wird auf der CPU ausgef hrt die Kernelfunktionen aus einer Queue startet die ihrerseits auf dem FPGA ausgef hrt werden 3 4 4 Architektur bergreifende C Sprachen f r den FPGA 3 4 4 1 OpenCL In der Spezifikation von OpenCL 36 wird angedeutet dass die Sprache f r alle m glichen Ar chitekturen unabh ngig vom Betriebssystem geeignet ist Dazu z hlen Vielkern CPUs GPUs Cell Prozessoren und DSP Weiter gibt es f r den eingebetteten Bereich eine Spezifikation die geringere Anforderungen an die Funktionalit t stellt Die folgende Tabelle listet existierende und verf
152. perator InstComp add Addition sub Subtraktion mul Multiplikation udiv vorzeichenlose Division sdiv Division urem vorzeichenlose Modulo srem Modulo shl links Schieben lshr rechts Schieben MSB aufgef llt mit Nullen and Und or Oder xor exklusives Oder InstGetElemPtr getelementptr Operator zur Adressberechnung von Vektoren und Matrizen InstSelect select bedingte Zuweisung InstTransfer load lesender Speicherzugriff store schreibender Speicherzugriff Tabelle 5 1 Abgeleitete Instruktions Klassen die LLVM Befehle repr sentieren und kurze Beschreibung Im LLVM Language Reference Manual 49 sind alle LLVM IR Befehle erl utert Tabelle 5 1 listet die Instruction Spezialisierungen auf und zeigt welche LLVM IR Befehle sie abdecken 74 5 3 VHDL Kompilierer k nnen i y Funktionsparameter gt Register Lesepfad k Y Schreibpfad Bild 5 7 Beispiel des generierten SSA AST der Kernelfunktion matrixada Als Beisliel der SSA AST Generierung wird wieder die Kernelfunktion matrixada aufgegrif fen Aus dem optimierten LLVM IR Quelltext entsteht das Speicherabbild des Syntaxbaums der mit dem Bild 5 7 veranschaulicht wird 5 OpenCL zu FPGA bersetzer 5 3 4 Parsebaum Analyse Der sourceAnalyser ben tigt ein sourceParser Objekt das bereits den SSA AST generiert hat Bild 5 8 zeigt dass nur die Methode parse ffentlich ist Sie ruft intern die privaten Metho den in entsprechender Reihenfo
153. r REG gelesen wird und wie viele 92 5 4 Rahmendesign ctrl_rd_reg_rdy clkPCle ram_wr_en clkUser Schreib ram_wr_addr rst RAM ram_wr_data i FIFO ram_rd_addr pcie_wr_data pcie_wr_addr ram_rd_en pole_wr_en O ram_rd_burst pcie_rd_addr o y idx_wr_id pcie_rd_addr_en Ro Schreib idx_wr_en pcie_rd_addr_size Zustandsmaschine ES idx_wr_data pcie_rd_data_get idx_rd_id ctrl_wr_reg ram_rd_data ctrl_wr_reg_en ram_rd_valid wr_reg_ idx_rd_data base_addr base_pos ctrl_rd_reg BEN pcie_rd_empty pcie_rd_data pcie_rd_count Bild 5 20 Schematisches Blockschaltbild der PCIe Einheit Worte die Lesel nge umfasst Danach ist die Zustandsmaschine in der Lage nachein ander mehrere Leseoperationen auszuf hren bis die Lesel nge ausgesch pft ist Die Lesedaten werden von den Schnittstellen geliefert und von der Logik gez hlt Ist die entsprechende Lesel nge an Worten an die DMA Schnittstelle gesendet worden ist die Zustandsmaschine f r die n chste Lesetransaktion bereit Der PCle Benutzer Adressraum ist auf die Schnittstellen RAM Index idx Register ctrl und Base aufgeteilt um vom Host aus auf diese zugreifen zu k nnen Die ram Schnittstelle hat 256 Bit breite Datenworte Aus diesem Grund existiert im Le sepfad ein separates FIFO f r den Speicher und die Register Das Lese RAM FIFO setzt die Wortbreite auf 32 Bit breite Datenworte f r die pcie Schnittstelle um Die idx Schnittstelle ist Teil de
154. r Register Schnittstelle auf der die Parallelindizes in Re gister der Kontrolleinheit geschrieben und gelesen werden k nnen Die base Schnittstelle ist Teil der Register Schnittstelle Sie existiert um die 26 Bit brei 93 5 OpenCL zu FPGA bersetzer ten Basisadressen in die Transferbl cke zu schreiben da diese erst zur Laufzeit bekannt werden Es existiert kein Lesepfad f r die Basisadressen e Die ctrl Schnittstelle ist Teil der Register Schnittstelle Sie f hrt zu einem 32 Bit breiten Kontrollregister in der Kontrolleinheit mit einem Lese und einem Schreibpfad 5 4 4 Datenflusskonzept In der Pipeline gibt es VHDL Transfer Bl cke die als Dateneinspeisepunkt und als Datenaus gabepunkt dienen Das Konzept einen Datenfluss in der Pipeline aufrecht zu erhalten sieht vor ein Bussystem mit dem Speicherkontroller zu verbinden ber das die Transferbl cke Zu griff zum Speicher bekommen Wenn in der Pipeline an einem Punkt keine Daten bereitstehen muss sie angehalten werden bis wieder Daten geliefert wurden Um die Pipeline am laufen zu halten ist eine kontinuierliche Versorgung mit Daten unerl sslich Aus diesem Grund besitzt jeder Transferblock einen Datenpuffer vergleichbar mit kleinen Caches Bei jedem Speicher zugriff wird ein Burstzugriff bestehend aus 128 Werten initiiert um m glichst effizient viele Werte zu lesen und in den Caches abzulegen Ein anderer Ansatz w re einen gro en Cache an den Date
155. r ein Thread pro Stream Multiprozessor laufen siehe Kapitel 2 1 2 Zur Bestimmung der effizientesten Implementierung wurde die Kernelfunktion IF siehe Ta belle 4 2 mehrfach in CUDA implementiert In der ersten Implementierung wurde der Zwi schenergebnisvektor unterteilt so dass jeder Abschnitt in den Shared Speicher passt Eine wei tere Implementierung lagert den Zwischenergebnisvektor in den gr eren lokalen Speicher aus Die optimale Implementierung liest blockweise eine Matrixzeile berechnet die Indizierung des Zwischenergebnisvektors und summiert das entsprechende Element Die Blockgr e ist gleich der Anzahl vorhandener Threads d h ein Block wird gleichzeitig gelesen und verarbeitet Die Index Berechnung basiert auf der Indizierung der gelesenen Daten i und j zu denen die Grau werte J und J aus der Lookup Tabelle bestimmt werden Die Summe aus und J indiziert das Element des Zwischenergebnisvektors auf den die zuvor gelesenen Daten addiert werden 4 3 2 8 Zwischenergebnisvektor Px y k Die Gr e des Zwischenergebnisvektors P x y k betr gt nur 4096 Elemente mitk 0 1 2 Ng 2 und Ng 4096 Auch dieser Zwischenergebnisvektor ist zu gro er w rde den Shared Spei cher komplett ausf llen und aufgrund des mangelnden Speichers keine weiteren Threads zur Ausf hrung zulassen F r die Implementierung ist die gleiche Strategie optimal wie sie auch beim Zwischenergeb nisvektor P y genutzt wird Der Unterschied liegt
156. r zweiten Ebene aus denen die GPC bestehen hei en Stream Multiprozessoren SM die sp ter weiter erl utert werden Erst in der dritten Ebene sind die Threadprozessoren TP CUDA Kerne zu finden Die drei stufige Struktur zu sehen in Bild 2 1 ist in der Lage bis zu 512 Threadprozessoren effizient zu verwenden Die schnellste in dieser Arbeit benutzte Grafikkarte NVIDIA GTX 480 mit der GF100 2 Grundlagen GPU Chip i SM Ebene d TP Ebene Bild 2 1 Hierarchieebenen der GF100 Architektur JT Architektur besteht aus vier GPCs Jeder GPC kann aus bis zu vier SMs bestehen Insgesamt gibt es in der GTX 480 allerdings nur 15 SMs aus den vier GPCs mal vier SMs wird ein SM bestimmt und ausgeschaltet Das wird gemacht um eventuelle Herstellungsfehler die den Aus fall einer SM zur Folge h tte miteinzubeziehen Entsprechend werden bei den anderen Gra fikkarten der GF100 Serie GTX 470 und GTX 465 mit Herstellungsfehlern h herer Dichte mehrere Einheiten abgestellt um dennoch eine funktionst chtige GPU mit geminderter Leis tung zu erhalten In jedem SM sind 32 TPs verbaut Zusammen besitzt die GTX480 eben 480 Threadprozessoren 15SM 32T P Auf der Chip Ebene sind bis zu sechs 64 Bit breite GDDRS5 Speicherkontroller ber ein Kom munikationsnetz an die GPC angeschlossen Die Gesamtbreite der Speicherschnittstelle um f
157. rallelisieren und als Quelltext in Kernel Funktionen abbilden Die Kernel Funktionen werden auf der GPU parallel ausgef hrt Die Skalierung der Parallelit t gemeint ist die Anzahl der Threads die die Kernel Funktionen durchlaufen sollen wird beim Funktionsaufruf mit der Angabe einer Dimension f r ein Grid und CUDA Bl cken definiert siehe Quelltext Beispiel 2 1 auf Seite 13 Das hei t innerhalb ei 10 2 1 Grafikkarten als Rechenbeschleuniger nes Grids werden CUDA Bl cke bestimmter Anzahl geschaffen die wiederum eine bestimmte Anzahl an Threads beinhalten Bild 2 3 veranschaulicht den Zusammenhang zwischen Grid Bl cken und Threads mit der M glichkeit einer mehrdimensionalen Anordnung Jeder Thread innerhalb der Bl cke f hrt die gleiche Kernelfunktion f r unterschiedliche Daten aus hnlich dem SIMD Prinzip nennt NVIDIA dieses Prinzip Single Instruktion Multiple Threads SIMT erw hnt in der CUDA Programmieranleitung 65 Thread 2d Block Feld Z 3d Thread Feld x x Bild 2 3 Anordnung vieler Threads in einem Grid und Bl cken Ein CUDA Grid besteht aus einem bis zu zweidimensional gro en Block Feld Ein CUDA Block kann eine bis zu dreidimensionale Anordnung von Threads sein Die Anzahl der Threads die beim Kernelaufruf bestimmt und angelegt werden werden von der Hardware dynamisch zur Laufzeit auf die Threadprozessoren verteilt Genauer wird je der CUDA Block vergleichbar wie ein Prozess auf einem Stream Mult
158. rdnungsgem e Funktion einzelner Teile zu ermitteln BPI Linear Flash 32 MB Der Flash Speicher h lt das Rahmendesign das mit dem Ein schalten des PCs und der Versorgung der Karte automatisch geladen wird 5 5 1 2 PCle Treiber und MPRACE Bibliothek Die Entwicklungssteckkarte ML605 als Beschleunigerkarte kommuniziert mit dem Host ber den PCle Bus Die Kommunikation wird mit Hilfe eines Linux Treibers PCle Treiber und einer C C Bibliothek MPRACE bew ltigt siehe 56 Der Puffermanager 57 aus der MPRACE Bibliothek der Transfers zwischen dem virtuellen und dem physikalischen Adress raum koordiniert arbeitet mit der bereits beschriebene DMA Engine zusammen Die MPRACE Bibliothek umfasst einfache Funktionen f r einen peripheren Transfer PIO einen direkten Speichertransfer DMA und Registerzugriffe in der DMA Engine die die Kommunikation steuert Die Kommunikation des Kompilierers basiert auf DMA Zugriffen und Registerzugrif fen Quelltextabschnitt 5 1 zeigt die verwendeten Lese und Schreibfunktionen virtual void readDMA const unsigned int address FPGA Adresse DMABuffer amp buf DMA Puffer const unsigned int count Anzahl zu lesender 32Bit Werte 1 2 3 4 5 const unsigned int offset 0 Versatz des ersten Wertes 6 const bool inc true Adresse im FPGA inkrementieren 7 const bool lock true Warte bis der Transfer fertig ist 8 const float timeout 0 0 Abbruchzeit in Millisekun
159. re Prozessoren abgebildet Der serialisierte Quelltext wird durch mehrfache Funktionsaufrufe wieder paralle lisiert die nacheinander auf den Softkernprozessoren ausgef hrt werden 3 4 3 2 OpenRCL OpenRCL ist eine Entwicklung die OpenCL f r den FPGA nutzbar macht Der Konferenz beitrag 52 von 2010 beschreibt die Architektur des entsprechenden FPGA Designs erl utert den bersetzungsvorgang und zeigt vergleichbare Ausf hrungsgeschwindigkeiten gegen ber GPUs Ziel der Entwicklung ist es erstens FPGAs f r die Architektur bergreifende Sprache OpenCL nutzbar zu machen und zweitens einen besseren Wirkungsgrad im Energieverbrauch zu erreichen als die Vergleichsarchitekturen GPU und CPU FPGA Design Bild 3 6 zeigt die unterschiedlichen Speicher ein Kommunikationsnetzwerk Kreuzschienenverteiler und die Threadprozessoren Jeder Threadprozessor PE hat einen privaten und einen lokalen Speicher P L Auf den privaten Speicher kann nur der anlie gende Threadprozessor zugreifen Auf den lokalen Speicher haben alle Threads der selben work group gegenseitigen Zugriff F r den Datenaustausch nutzen sie den Kreuzschienenver 41 3 Stand der Technik teiler mit eignen Speichern SP der mit Blick auf eine maximale Leitungsl nge im FPGA sehr viele Kommunikationsteilnehmer erlaubt Auf den global geteilten Speicher k nnen al le Threads direkt zugreifen Die Architektur ist der einer GPU sehr hnlich Tats chlich sind Kreuzschi
160. rer f r OpenCL zu entwickeln als f r CUDA da OpenCL einem offenen Standard unterliegt Es gibt einen standardisierten und gut dokumentierten Funktionssatz was bei CUDA nicht der Fall ist e Weiter ist OpenCL oder CUDA eine weit verbreitete Sprache die von vielen Pro grammierern akzeptiert wurde Das macht eine Einarbeitung in die Sprache n tzlich da es mehrere Anwendungsgebiete gibt Sprachen die mit Sprachkonstrukten erweitert werden m ssen haben es schwerer sich durchzusetzen 65 5 OpenCL zu FPGA bersetzer Die weiteren Gedanken lehnen daran an wie der Kompilierer OpenCL in ein FPGA Design umsetzt Die OpenRCL Entwicklung simuliert eine GPU hnliche Architektur im FPGA f r die ein Programm bersetzt werden kann Der Ansatz dieser Arbeit sieht vor einen Pipeline generator zu entwickeln der das Programm in eine Hardwarepipeline bersetzt Somit werden die Ressourcen zielgerichtet f r das OpenCL Programm eingesetzt und das Pipelinekonzept verspricht eine hohe und effiziente Auslastung der synthetisieren Hardware 5 2 bersicht Die entwickelten Bestandteile des bersetzer Konzepts sind miteinander verwoben Aus die sem Grund wird hier eine bersicht in Bild 5 1 gegeben die nur im Groben die Bestandteile beschreibt und deren Interaktionen mit Pfeilen darlegt OpenCL main OpenCL kernel OpenCL Laufzeitumgebung VHDL Kompilierer Rahmendesign 3 Pipeline Bild 5 1 B
161. reset Die Informationsbits geben Auskunft ber einen FIFO berlauf und die Pipeline ID e Eine Zustandsmaschine verwaltet die Berechnung Sie verharrt in einem Ruhezustand in dem die Pipeline nicht rechnet bis das nullte Bit des Kontrollregisters gesetzt wird Der Berechnungsstart wird der Speicherverwaltungseinheit und der Pipeline mitgeteilt Im n chsten Zustand wird gewartet bis die Pipeline das Ende der Berechnung f r die aktu elle Kombination der Parallelindizes meldet Dann wird die Pipeline zur ckgesetzt und gepr ft ob alle Parallelindizes Kombinationen ausgesch pft wurden An dieser Stelle verzweigt sich die Zustandsmaschine entweder startet die Berechnung mit neuen Paral lelindizes Kombinationen oder alle Berechnungen sind fertig und die Zustandsmaschine geht in ihren Ruhezustand zur ck Das nullte Bit des Kontrollwortes wird zur ckgesetzt e Die Parallelindizes Kombination wird zwischen jeder Berechnugnsphase ermittelt Wel che gerade berechnet wird bestimmen die Register work group id x bis z In den Be rechnungspausen wird die work group id erh ht wie bei einer dreifach verschachtelten Schleife die ein Volumen elementweise durchl uft Die work group size entspricht der Schleifengr e von jeder Dimension mit der sich die Anzahl aller Pipelineberechnun gen ermitteln l sst Die Pipeline berechnet alle Threads einer work group Die Anzahl der Threads l sst sich von den Registern work item size x bis z bestimmen Da
162. ringer als die einer Vielkernrechnerarchitektur e Der gravierende Vorteil einer OpenCL Implementierung ist die doppelte Einsatzf hig keit In dieser Arbeit k nnen Algorithmen die f r die Online Prozessierung FPGA entwickelt werden dann auch f r die Offline Prozessierung GPU eingesetzt werden Der Ansatz dieser Arbeit einen OpenCL Kompilierer zu entwickeln ist vergleichbar mit der Mischung der drei gezeigter Ans tze Erstens es wird eine hnliche Pipelinegenerierung wie die von CHiMPS verwendet zweitens wird die parallele Sprache OpenCL f r den FPGA be nutzt nach der Grundidee von OpenRCL und drittens lehnt der bersetzungsvorgang an Tri dent an LLVM verwenden zu wollen Die Mischung der drei existierenden Ans tze zu einem besitzt Vorteile gegen ber jedem der einzeln gezeigten Dieser Ansatz vereint die Vorteile e FPGAs mit der einfachen parallelen Programmiersprache OpenCL beschreiben zu k n nen 46 3 4 Software Hardware Kompilierer e ber einen Pipelinegenerator zu verf gen der die Vorteile der Streambarkeit und des hohen Rechendurchsatzes nutzt e Die bersetzung auf LLVM aufzubauen um die Optimierer nutzen zu k nnen und die bersetzung einfach entwickeln zu k nnen 47 4 Haralick Algorithmus GPU beschleunigt 4 1 Untersuchung des Haralick Algorithmusses Der Haralick Algorithmus besteht aus zwei Teilen Im ersten Teil werden von den Mikroskop bildern Co occurrence Matrizen erstel
163. rst ndlicherweise liefern nicht alle aufgenomme nen Daten n tzliche Ergebnisse Mit dem beschleunigten Mikroskop ndert sich die Situation dahingehend dass nicht nur weniger Zeit f r die Datenaufnahme ben tigt wird sondern auch dass der Datendurchsatz und die anfallende zu speichernde Datenmenge um den Beschleuni gungsfaktor 20 ansteigen Auch die Diskrepanz zwischen Aufnahmezeit und Datenanalysezeit verschlechtert sich weiter und mehr kostenintensive Daten Cluster werden ben tigt die viel Energie verbrauchen um die Daten abrufbereit zu halten Eine Beschleunigung der Offline Prozessierung ist au erordentlich wichtig um e die Datenmengen die gespeichert werden m ssen zu reduzieren e rechenintensive Algorithmen in der Offline Prozessierung verwenden zu k nnen e den Geschwindigkeitsgewinn des beschleunigten Mikroskops auf die Offline Prozes sierung zu bertragen um Ergebnisse schneller auswerten zu k nnen Eine Konsequenz ist den bestehenden Rechen Cluster durch weitere Rechner CPUs zu er weitern CPUs eignen sich f r alle Rechenaufgaben was bedeutet dass sich durch eine ver mehrte Anzahl eine Beschleunigung erzielen l sst sofern die Anzahl nicht zu gro wird Eben so kann eine Beschleunigung durch Beschleunigerkarten realisiert werden die in den bestehen den Rechner Knoten eingesetzt werden Beispiele f r Beschleunigerkarten sind FPGAs und GPUs In der Regel wird von ihnen eine h here Beschleunigung gegen b
164. rt als Ergebnis bernommen In dieser Implementierung darf der boolesche Ausdruck lediglich eine Vergleichsoperation mit einer Konstanten sein F r eine Demonstration gen gt die Annahme einer Konstan te wodurch bereits viele Beispielprogramme bersetzungsf hig werden Die Konstante l sst sich in einer Weiterentwicklung durch eine Variable ersetzen indem der Parser an gepasst wird analysePipelinePart UNd analyseTransferPart weisen alle arithmetischen Operatoren der Pipeline oder des Datentransfers zu Die Pipeline soll nur die Operationen enthalten die zur Ergebnisberechnung beitragen Die Operationen die zur Berechnung der Adresse ben tigt wer den beispielsweise die Indexberechnung eines Elements in einem Array Index Basis i soll in einer speziellen Komponente des Pipeline Baukastens passieren Wie die beiden Analyse methoden die Zuordnung herausfinden wird am Beispiel einer Matrixaddition verdeutlicht Die Kernelfunktion besteht im wesentlichen aus der Zeile Clidy w idx Alidy w idx Blidy x w idx Der daraus entstehende SSA AST ist in Bild 5 9 dargestellt 77 5 OpenCL zu FPGA bersetzer InstCall global_id 1 ta i InstComp InstCall mul global_id 0 aa add d N getelementptr getelementptr J 4 InstTransfer InstTransfer load load getelementptr Pipelineoperation Adressoperation InstTransfer store Funktionsparameter
165. ruktur bei der eine Untermenge der Quellzeile eingelesen wird und eine Matrix Block Struktur bei der ein Ausschnitt aus der Quellmatrix eingelesen wird Aufgrund eines besseren Verhaltens beim Datentransport wird f r die Bildmerkmale f5 f6 und f3 die Matrix Block Struktur zur Summenbildung verwendet 60 4 3 GPU Implementierung 4 3 2 6 Index abh ngige Merkmal Gleichungen Ein Blick auf die Gleichungen der Bildmerkmale und der Definitionen zeigt dass einige den Index mit in die Berechnung einbeziehen Gemeint ist dass der Grauwert entsprechend der Po sition in der vollen Co Matrix mit einbezogen wird Indexabh ngige Gleichungen sind 4 3 4 4 4 5 4 14 4 15 4 17 und 4 18 die die Bildmerkmale f3 f4 f5 und die De finitionen Py y Py mean und var berechnen Aufgrund der gepackten Co Matrix entsteht f r die Berechnung der indexabh ngigen Gleichungen ein Mehraufwand da ein zus tzlicher Lesezugriff auch die Lookup Tabellen f r die Grauwertbestimmung notwendig ist Um den Geschwindigkeitseinbruch klein zu halten wird ein Texturencache f r die Lesezugriffe auf die Lookup Tabelle verwendet 4 3 2 7 Zwischenergebnisvektor Px y k Der Zwischenergebnisvektor P k mit k 2 3 2Ng 2 und Ng 4096 hat eine Gr e von 8188 Elementen zu gro f r den Shared Speicher mit 4096 Elementen 16kBytes 4By tes pro Element Falls doch der gesamte Shared Speicher verwendet wird k nnte aufgrund der Knappheit nu
166. s eine Vorschrift angewendet wie ein entsprechender Knoten im Block AST zu generieren ist Dieser Vorgang ist in private Hilfsmethoden unterteilt worden die von compile aufgerufen werden Im privaten Bereich unter der Instanzvariable vnaiast wird der Block AST entste hen Die Hilfsmethode findBlockInast sucht im Block AST eine bestimmte Stelle an der 79 5 OpenCL zu FPGA bersetzer die Hilfsmethode insertBlockast einen neuen Block einf gen kann Die Knoten des Block ASTs sind Spezialisierungen der abstrakten Klasse vnaiBlock Die Ab geleiteten Klassen werden im n chsten Abschnitt vorgestellt 5 3 5 2 Zuordnung Instruktion zum VHDL Block Der Pipeline Baukasten enth lt f nf VHDL Bl cke aus denen die Pipeline zusammengesetzt wird Mit wenigen Bl cken auszukommen hat den Vorteil dass ein Zusammensetzen wenig komplex ist D h es gibt nicht viele M glichkeiten die Bl cke miteinander zu verbinden Der entstehende Nachteil ist dass viel Funktionalit t aus der OpenCL Kernelfunktion auf wenig Bl cke verteilt werden muss Diese Arbeit verfolgt den Ansatz die Funktionalit t der VHDL Bl cke f r Standardf lle auszulegen und zu optimieren Die Tabelle 5 2 geht in der ersten Spalte von den Instruktion Klasseninstanzen aus die in der zweiten Spalte auf VHDL Bl cke widergespiegelt werden Klassenname VHDL Block Kommentar InstCall Wird nicht in einen VHDL Block bersetzt InstCmp BlockCondAss
167. schleunigung durch den Einsatz von Co Prozessoren zu erwarten ist l sst sich mit der Formel von Amdahl 3 1 berechnen Der Parameter n steht f r die m gliche Parallelisierung d h auf wie viele Prozessoren die Rechenzeit aufgeteilt wird Die Formel hat den Ursprung in Amdahls Ver ffentlichung von 1967 20 Besser und mit vielen Beispielen erl utert ist sie in 40 Computer Architecture Beschleunigung 3 1 ts 2 Mit genauer Betrachtung der Formel wird klar dass f r eine gro e Beschleunigung die Aus f hrungszeit des seriellen Programmteils deutlich kleiner sein muss als die des parallelen Pro grammteils Die theoretisch maximale Beschleunigung konvergiert schnell gegen einen maxi malen Wert Intuitiv kann bei einem 50 prozentigen parallelen Programmteil eine maximale Beschleunigung von zwei erwartet werden bei 90 Prozent ergibt das einen Faktor von Zehn Ob sich der Aufwand einer Beschleunigung lohnt kann somit genau abgewogen werden Daher ist es wichtig die genauen Ausf hrungszeiten der seriellen und der parallelen Programmteile zu kennen Profiling Werkzeuge sind in der Lage ein Profil der Ausf hrungszeiten aller Funk tionen zu ermitteln Wei man welcher Teil der Anwendung am langsamsten ist geht es darum diesen zu be schleunigen Der erste Ansatz sollte sein den Algorithmus mathematisch zu untersuchen Gibt es die M glichkeit die Formeln zu vereinfachen um Rechenoperationen einzusparen oder re cheninten
168. se Schreibtransfers auf den Speicher Es gibt Anforderungen an die Basisadresse und die Zugriffsblockgr fe beim Lesen bzw beim Schreiben auf den globalen Speicher Die kleins te Zugriffsgr e besteht aus einem zusammenh ngenden Block von 64 Bytes der mit einem 102 5 5 OpenCL Laufzeitumgebung Vielfachen von 64 Bytes und der genauen Bytegrenze von allen 64 Bytes adressiert werden muss Der DMA Transfer besitzt die Anforderung einer konfigurierbaren maximalen Daten l nge Die Methoden readram und writeram ber cksichtigen alle drei Anforderungen und setzen beliebige Adressen und Datenl ngen wenn n tig in mehrere DMA Transfers um Reset Funktionen Vor jedem DMA Transfer muss die entsprechende Lese bzw Schreib DMA Engine r ckgesetzt und die Abbruchzeiten neu gesetzt werden Die Methode resetPCIe o bernimmt diese Aufgabe Sie wird von den Transferfunktionen vor jedem DMA Transfer aufgerufen Die Methode resetPipeline setzt ein Bit im Kontrollregister das f r wenige Taktzyklen ein Reset in der Pipeline durchf hrt Lese Schreibtransfers auf die Register Beim Lesen und Schreiben der Register in der Kon trolleinheit des Rahmendesigns implizieren die Methoden readreg und writereg die Re gisteradresse im Adressraum so dass zur Adressierung lediglich die gew nschte Registernum mer angegeben werden muss Gr e der parallelen Indizes setzen Mit den Methoden writeworkGroupSize und writework ItemSize
169. sign synthetisiert pipelineTCL sh build pipeline vha Das Skript generiert lediglich die Hardware Pipeline als dynamischen Teil und wird zum bersetzen statischen Rahmendesigns verkn pft Der n chste Schritt ist die Pipeline auf dem Chip auszutauschen Weitere Details sind im Abschnitt Pipeline Module mit DPR austau schen 5 5 3 2 erl utert 5 5 2 4 Daten bertragung und Pipeline starten Befindet sich die Kernelfunktion als Pipeline im FPGA Kann sie f r eine Berechnung heran gezogen werden Der Quelltextausschnitt 5 4 zeigt welche Strukturen ben tigt werden um die Berechnung auszuf hren kernel clCreateKernel program kernel amp status queue clCreateCommandQueue context device 0 amp status meml clCreateBuffer context CL_MEM_READ_WRITE sizeof unsigned int MEGA NULL amp status mem2 clCreateBuffer context CL_MEM_READ_WRITE sizeof unsigned int MEGA NULL amp status h_meml new unsigned int MEGA h_mem2 new unsigned int MEGA status clEnqueueWriteBuffer queue meml true 0 sizeof unsigned int x MEGA h_meml 0 NULL NULL status clSetKernelArg kernel 0 sizeof meml meml status clSetKernelArg kernel 1 sizeof meml meml status clSetKernelArg kernel 2 sizeof mem2 mem2 107 14 15 16 18 5 OpenCL zu FPGA bersetzer const size_t global_work_size 3 MEGA 1 1 const size_t local_work_size 3 1 1 1 status
170. sive Operationen Divisionen Logarithmus zu vermeiden bietet das die gr t 23 3 Stand der Technik m gliche Beschleunigung Schneller als gar nicht Rechnen zu m ssen geht nicht Wenn der Algorithmus seriell ist kann man in der Literatur nach einer parallelen Version suchen Es kann durchaus sein dass der Algorithmus f r die bisher seriell rechnende Welt serialisiert wurde und urspr nglich parallel war Erst der letzte Schritt ist die Portierung und Anpassung an die Co Prozessor Architektur beispielsweise auf eine GPU oder einen FPGA In erster Linie wird durch eine Parallelisierung beschleunigt Die einfachste Variante ist die Parallelisierung der Anwendung nach Flynn MISD 28 indem die Anwendung mehrfach ausgef hrt wird Die Parallelisierung des Datenflusses SIMD betrifft die Anwendung selbst in dem die Operatoren gleichzeitig auf mehreren Werte angesetzt werden SIMD Eine Paralle lisierung mit einer Pipeline basiert darauf das ganze Programm in Funktionseinheiten gleicher L nge zu unterteilen In jedem Takt k nnen die Funktionseinheiten ihr Teilprogramm berech nen und die Ergebnisse an die n chste Einheit weiterreichen Dieses Prinzip wird bei der Pro zessorentwicklung genutzt die Taktrate einer Abarbeitung von Instruktionen zu steigern Bei einer Parallelisierung eines Programms liegt der Gewinn in der gleichzeitigen Ausf hrung al ler berechnenden Operatoren im Programm zu jedem Takt 3 1 2 GPGPU
171. ss die Anwei sungen an geeigneten Stellen im Quelltext erg nzen Cache Anweisungen Der Programmierer kann Cachegr e Zeilenl nge Anzahl der B nke Cache Lese Schreibrechte konfigurieren und hat damit einen entscheidenden Einfluss auf den optimalen Datenfluss Separate Speicher Das Einbinden weiterer Speicher reduziert den Zugriff auf den langsame ren Hauptspeicher Implementierungsstil Der FPGA kann eine TCL Operation mit unterschiedlichen Ressour cen umsetzen Mit Anweisungen wird dem Kompilierer mitgeteilt welche Operation beispiels weise mit Logikzellen mit DSP Slices oder sogar mit einem Softcore Prozessor implementiert wird Schleifen aufrollen Die Anzahl der Schleifendurchl ufe kann durch Daten Parallelit t redu ziert oder gar ganz aufgel st werden Integer Bitbreiten spezifizieren Nicht in jedem Programm wird die Bitgenauigkeit eines In tegers voll ausgenutzt Kennt der Programmierer den Wertebereich der Berechnungen k nnen durch Eingrenzungen der Bitbreite FPGA Ressourcen eingespart werden 3 4 3 Parallele C Sprachen f r den FPGA 3 4 3 1 FCUDA FCUDA 67 ist eine CUDA Erweiterung die Kernels auf den FPGA abbildet Mit erg nzen den Anweisungen werden die Kernel Funktionen in eine C Variante namens AutoPilot C 86 bersetzt Die Syntheseprogramme von AutoPilot erstellen eine Netzliste die mit der Xilinx Werkzeugkette das Design auf den FPGA zur Ausf hrung bringt Der erste bersetzungsschritt in Bil
172. sse in K stchen Alle Bl tter gekennzeichnet mit ei nem Stern sind ihrerseits von den Co Matrizen abh ngig Der Abh ngigkeitsgraph bietet mehrere Optimierungen f r eine Implementierung Wie bereits erw hnt zeigt er die optimale Reihenfolge der Berechnungen um doppelte zu vermeiden Des weiteren k nnen Programmstrukturen abgeleitet werden Alle Bl tter eines Zweigs k nnen ge meinsam in einer Funktion bzw Schleife implementiert werden Daraus ergibt sich dass f r diesen Zweig die Quellen nur einmalig gelesen werden was wiederum unn tige Speichertrans fers reduziert Ebenso wird das Speicherzugiffsverhalten auf kleinere Regionen des Speichers konzentriert worin Vorteile f r Architekturen mit Caches entstehen 54 4 2 CPU Implementierung 4 2 CPU Implementierung Das erste Ziel war die bereits existierende Software Version zu optimieren um sie auf einem Computercluster als Ein CPU Applikation laufen zu lassen Durch mehrfache Ausf hrung mit unterschiedlichen Zellbildern kann der gesamte Rechenaufwand parallelisiert werden Bild 4 4 zeigt die Softwarestruktur in einem Struktogramm Die u ere Schleife iteriert ber alle Zellen C eines Multizellbildes Innerhalb der Schleife werden alle Co Matrizen generiert Weiter folgt eine Doppelschleife ber den Winkel A und Distanz D Das hei t f r jede Zel le C existieren A D Co Matrizen die f r alle 13 Bildmerkmale seriell berechnet werden
173. sse vom Grafikkartenspeicher in den Hauptspeicher des PCs kopiert Bild 2 1 zeigt den Quelltext einer Kernelfunktion die ei ner Matrixaddition berechnet und somit eine zweidimensionale Threadanordnung des Grids und der Bl cke nutzt __Gglobal__ void CUDA_matrixAdd float A float B float C Anzahl Threads in einem Block int blockSize blockDim y x blockDim x aktuell ausgef hrte Blocknummer int blockNum blockIdx y x gridDim x blockIdx x aktuell ausgef hrte Threadnummer im Block int threadNum threadldx y blockDim x threadIdx x globale Threadnummer von allen Bl cken int globalNum blockNum blockSize threadNum Jeder Thread adressiert zwei Speicherzellen im Hauptspeicher addiert und schreibt sie zurueck C globalNum A globalNum B globalNum int main Kernelaufruf mit ca 50000 Threads CUDA_matrixAdd lt lt lt 256 192 gt gt gt A B C Quelltext 2 1 Von vielen Threads durchlaufene CUDA Kernelfunktion zur Berechnung einer Matrixaddition Die in der Kernelfunktion eingebauten Variablen gridDim blockDim blockIdx und threadIdx existieren ohne Deklaration und sind Strukturen mit jeweils drei Komponenten x y und z Sie zeigen an welcher Thread innerhalb welchen Blocks gerade in der Ausf hrung ist Im Quelltextbeispiel wird in lokalen Variablen die Blockgr e die Blocknummer im Grid und die Threadnummer im Block berechnet Die Blocknummer und Threadnummer berechn
174. t FPGA Beschleunigerkarte und dem externen Speicher ist sehr komplex im Vergleich zur eben gezeigten Vektor Additions Funktion F r die Aufgaben einen Datenfluss herzustellen ben tigen Programmierer sehr viel Zeit und umfang reiche Hardwarekenntnisse Die OpenCL Implementierung bietet eine Umgebung OpenCL Laufzeitumgebung in der die Kommunikation mit der FPGA Beschleunigerkarte und die Ge nerierung eines Datenflusses bereits implementiert ist Der n tige Progammaquelltext der die Laufzeitumgebung verwendet und die Kommunikation mit der Beschleunigerkarte regelt ist f r alle OpenCL Implementierungen sehr hnlich weit verbreitet und leicht zu erlernen Wei ter bietet die OpenCL Implementierung einen bersetzer an OpenCL zu FPGA Kompilierer eine C hnliche parallele Funktion in eine Hardwarebeschreibungssprache zu kompilieren Mit der Laufzeitumgebung und dem bersetzer erm glicht diese OpenCL Implementierung Programmierern ohne Hardwarekenntnisse eine FPGA Beschleunigerkarte einzusetzen ber setzte Funktionen k nnen zur Laufzeit im FPGA geladen und zur Ausf hrung gebracht werden Die bersetzungszeit einer OpenCL Kernelfunktion in ein Pipelinemodul mit weniger als zehn Operatoren betr gt ca 40 Minuten Die bersetzungszeit mit der Datenfluss und Kontrolllogik PCIe Core DMA Engine Speicherkontroller und Rahmendesign ben tigt mehr Zeit Bereits bersetzte Pipelinemodule k nnen im Bruchteil einer Sekunde in den FPGA gela
175. t dass nur noch 278 Zeilen und Spalten gespeichert werden w hrend die Co Matrix aus Bild 4 1 a in voller Gr e 4096 Zeilen und Spalten ben tigt F r dieses Beispiel Konnte der Speicherplatzbedarf von 64Mbyte auf 300kBytes reduziert werden Im Durchschnitt liegt die gepackte Co Matrix Gr e bei 1 5 MBytes F r die sp tere Bildmerkmalsberechnung ist die Position eines Wertes in der Co Matrix von Bedeutung Die Position repr sentiert einen Grauwert in der vollen Co Matrix und flie t bei manchen Bildmerkmalen mit in die Berechnung ein Um sp ter den Grauwert bestimmen zu k nnen wird neben der gepackten Co Matrix eine Index Grauwert Tabelle konstruiert Kompressionsmethoden die sp rlich besetzte Matrizen als Position Werte Paar in einer Struk turliste speichern haben den Nachteil die Elemente indirekt adressieren zu m ssen was den Speicherzugriff erheblich verlangsamen w rde Diese Methode der gepackten Matrizen stellt einen Kompromiss zwischen wenig Speicherbedarf und direkten Speicherzugriffen dar 50 4 1 Untersuchung des Haralick Algorithmusses 4 1 2 Haralick Textur Merkmale Die Haralick Texturen Merkmale umfassen 14 Bildmerkmale zusammengefasst in 73 In die ser Implementierung wurden die Bildmerkmale 4 1 bis 4 13 optimiert Bildmerkmal Num mer 14 Maximul Correlation Coefficient wurde bereits in der Implementierung der Biologen weggelassen Sie haben festgestellt das Bildmerkmal Nummer 14 keinen Beitr
176. t die Vergleichsfunktion mit der die Daten der cmp Schnittstelle mit dem konstanten Ausdruck CMP_CONSTANT vergli chen werden Implementierte Vergleichsfunktionen sind gleich ungleich gr er gr er gleich kleiner und kleiner gleich e Das boolesche Vergleichsergebnis steuert einen Multiplexer an der entweder den Daten kanal a oder b auf den Datenausgang c weiterleitet 88 5 3 VHDL Kompilierer e S mtliche Dateneing nge sind mit Verz gerungsgliedern ausgestattet die sich mit den DELAY Parametern konfigurieren lassen 5 3 6 3 Zus tzliche Logik Dieser Abschnitt entspricht dem fehlenden bersetzungsschritt 3 aus dem Blockschaubild 5 4 F r die Pipeline ist zus tzliche Logik erforderlich die mit der bersetzung generiert werden muss Ohne die zus tzliche Logik k nnte mit der benachbarten Logik dem Rahmendesign keine ordnungsgem e Funktion entstehen Speicherbus Oder Logik In der Pipeline k nnen sich mehrere Transferbl cke befinden die alle an den Speicherbus angeschlossen sein m ssen Es gibt mehrere Varianten viele Quell signale zu einem zu kombinieren Eine Variante ist das Multiplexen bei dem ein Kontrollsi gnal bestimmt welches Quellsignal auf den Bus getrieben wird Eine andere Variante ist die Tristatelogik wobei alle Quellsignale elektrisch miteinander verbunden sind aber immer nur eines senden darf w hrend die anderen einen hochohmigen Zustand einnehme
177. t werden automatisch die Elemente 0 bis 31 gelesen somit werden mit einem Spei cherzugriff die Threads O bis 31 des gleichen Warps bedient Im Benutzerhandbuch wird dies als coalescing bezeichnet Eine weitere Praktik auf die geachtet werden muss betrifft wieder die Speichertrans fers Bei der Adressierung einer Matrix oder eines Volumens muss jede Basisadresse einer Zeile ein vielfaches der Warp Gr e sein M chte man die Elemente 33 34 aus dem Speicher lesen werden aber die Elemente 31 32 33 34 angefordert vielfaches der Warp Gr e In diesem Fall sind die Elemente 31 und 32 unn tigerweise gelesen worden Verhindert wird dies indem jede Zeile in der L nge erweitert wird damit die n chste Zeile eine g ltige Basisadresse besitzt Dabei kann am Ende jeder Zeile ein Da tenbereich entstehen der nicht verwendet wird Im Benutzerhandbuch mit dem Begriff pitch zu finden Der Shared Speicher ist in B nke unterteilt die unterschiedliche Speichermodule repr sentieren Wenn viele Threads auf den Shared Speicher gleichzeitig zugreifen wollen muss garantiert werden dass jeder Thread eine unterschiedliche Bank adressiert Im Fall eines Bank Konflikts wenn mehrere Threads auf die selbe Bank zugreifen werden die Transfers serialisiert Dies gilt es zu vermeiden Siehe im Benutzerhandbuch unter dem Begriff bank conflict Wie oben bereits beschrieben sollten alle Threads eines Warps dem selben Ausf hrungs pfad folgen Da eine
178. teraktionen zwischen Viren mit Zellen geforscht Bereits in der Vergangenheit wurden Zellen mit Viren infiziert und es wurde beobachtet wie die Viren in die Zelle eindringen sich vermehren und wieder austreten Die Untersuchung konzentrierte sich auf einzelne Proteine oder Gene was auch der Grund daf r ist nur wenig allgemeine Aussagen treffen zu k nnen Ein systembiologischer Ansatz verspricht allgemeinere und zielgerichtete Aussagen treffen zu k nnen Die Systembiologie ist eine junge Disziplin in der die Zelle als biologisches System mit allen vorhandenen Zellbausteinen biologischen Prozessen und Interaktionsm glichkeiten in einem mathematischen Modell betrachtet wird Lebensprozesse mit mathematischen For meln zu beschreiben ist keine leichte Aufgabe aber der Erkenntnisgewinn wird gro sein Wirkungsvolle Medikamente mit m glichst wenigen Nebenwirkungen herstellen zu k nnen ist ein Zukunftsziel dieser Forschung An einem vereinfachten Beispiel soll verdeutlicht werden wie eine systembiologische Anwen dung funktioniert Zellen besitzen nach heutigem Kenntnisstand ca 20 25 Tausend Gene in denen alle vererbba ren Informationen gespeichert sind Einzelne Gene oder auch Kombinationen davon sind f r die Proteinproduktion Enzyme verantwortlich Die Enzyme werden f r chemische Reaktio nen ben tigt beispielsweis einer Zellteilung aber auch Viren brauchen Enzyme als Nahrung 1 Einf hrung und Ziele des Viroquant Projekts
179. tragenden Protokollinformationen und einer nicht optimalen Ausnutzung aller Takte aufgrund von Latenzzeiten Die DMA Engine steht auf der Internetseite von OpenCores 29 bzw den Institutsseiten der Universit t Heidelberg 47 zum Herunterladen zur Verf gung Details der DMA Engine sind im CBM Bericht 30 nachzu lesen Aus Designsicht bietet der PCle Core zwei einfache Benutzerschnittstellen zum Lesen und Schreiben an die mit der PCIe Einheit im Rahmendesign verbunden sind 5 4 3 PCle Einheit In Bild 5 20 ist der funktionale Aufbau der PCIe Einheit aus dem Rahmendesign schematisch dargestellt e Die PCIe Einheit verf gt ber mehrere Dualport FIFOs die unterschiedlich getaktet sind und somit neben der Datenspeicherung auch die Aufgabe der Taktumsetzung haben Die Signale von und zur DMA Engine clkPCle sind mit 125 MHz getaktet die Signale von und zu den benachbarten Einheiten clkUser sind mit 133 MHz getaktet e Es gibt zwei DMA Schnittstellen Eine um Daten aus dem FPGA zum Host zu trans ferieren Lesen und eine die Daten vom Host zum FPGA transferiert Schreiben Die Signalnamen beider Schnittstellen im Bild beginnen mit dem Pr fix pcie e Beim Schreib DMA wird ber die Adresse das Ziel FIFO zum Speicher RAM oder zu den Registern REG bestimmt und die Daten werden dahin geleitet e Der DMA Lesevorgang wird von einer Zustandsmaschine gesteuert Sie bestimmt ber die Adresse von welcher Schnittstelle RAM ode
180. ts in FPGAs In Proceedings of the 2006 ACM SIGDA 14th international symposium on Field programmable gate arrays FPGA 06 Seite 12 20 ACM New York NY USA 2006 doi 10 1145 1117201 1117204 Pavle Belanovic und Miriam Leeser A Library of Parameterized Floating Point Modules and Their Use In Proceedings of the Reconfigurable Computing Is Going Mainstream Literatur 24 25 26 27 28 29 30 31 32 33 12th International Conference on Field Programmable Logic and Applications FPL 02 Seite 657 666 Springer Verlag London UK UK 2002 Ian Buck Tim Foley Daniel Horn Jeremy Sugerman Kayvon Fatahalian Mike Houston und Pat Hanrahan Brook for GPUs stream computing on graphics hardware ACM Trans Graph Volume 23 777 786 August 2004 doi 10 1145 1015706 1015800 C Conrad H Erfle P Warnat N Daigle T L rch J Ellenberg R Pepperkok und R Eils Automatic identification of subcellular phenotypes on human cell arrays Genome Res Volume 14 6 1130 6 2004 doi 10 1101 gr 2383804 Qualis Design Corporation VHDL Quick Reference Card URL http www eda org rassp vhdl guidelines vhdlgrc paf 1995 Ron Cytron Jeanne Ferrante Barry K Rosen Mark N Wegman und F Kenneth Zadeck Efficiently computing static single assignment form and the control dependence graph ACM Trans Program Lang Syst Volume 13 451 490 October 1991 doi 10 1145 115372 115320 M Flyn
181. ufe mit Register Beispielsweise liegt das Ergeb nis nach zwei Takten im Block 2 fest Der rechte Teil der Zuweisung entspricht einer kombinatorischen Logik und darf mehrere Operatoren enthalten wobei die maximale Durchlaufzeit die Taktrate mindert e Der Kanal 1 entspricht einem FiFo der zwischen den beiden seriellen Pfaden im pa rallelem Segment einen Datenaustausch zul sst Dabei verh lt sich das FiFo wie eine Synchronisation da normalerweise Block 2 gleichzeitig mit Block 4 ausgef hrt werden w rde Der Operator speichert einen Wert im Kanal ab der Operator wartet so lange bis ein Wert bereit liegt und liest ihn aus Handel C hat einige Einschr nkungen gegen ber ANSI C die f r eine Hardware bersetzung plausibel sind Es werden keinerlei Flie8kommadatentypen oder union Datentypen unterst tzt es gibt keine dynamische Speicherverwaltung mit malloc oder free und f r die Hardware schwierig realisierbare Rekursionen funktionieren nicht Es gibt weitere Sprachkonstrukte die von Handel C nicht bersetzt werden beispielsweise darf die sizeo Funktion oder die main Funktion keine Parameter und keinen R ckgabewert haben Weitere ANSI C Sprachelemente die in Handel C nicht konform sind lassen sich im Benutzerhandbuch nachschlagen 3 4 2 2 TRIDENT Trident 74 kompiliert C C Funktionen mit 1oat oder double Berechnungen in eine Hard warebeschreibungssprache und bietet die Synthese in ein FPGA Design Der Kompilierer gi
182. ung und Verdrahtung des gesamten FPGAs In der Regel spart dieser Vorgang Zeit und bei einer bersetzung vieler Partitionen k nnen die Xilinx Werk zeuge parallel auf Mehrkernprozessoren arbeiten Eine dynamische partielle Rekonfiguration DPR gemeint ist eine teilweise Neuprogrammie rung eines FPGAs ist der Austausch einer Partition mit einer anderen w hrend des Betriebs Voraussetzungen f r das Funktionieren sind im DPR Benutzerhandbuch 80 ausf hrlich er l utert und hier zusammengefasst e Die Schnittstellen der beiden Partitionen m ssen gleich sein e Der Austausch gelingt nur zwischen dynamischen Partitionen dynamisches Design und statischer Logik statisches Design e Es m ssen auf jeder Seite statisches und dynamisches Design vor der Schnittstelle Re gisterstufen sein Diese Anforderung existiert weil die Xilinx Tools ber diese Grenze das Zeitverhalten nicht optimieren k nnen e Vor dem Austausch muss die Grenze geschlossen werden Das bedeutet alle Signale vom dynamischen zum statischen Design m ssen logisch entkoppelt werden W hrend der Programmierung des dynamischen Designs k nnen dessen Ausg nge undefinierte Werte annehmen e Das gesamte Design ben tigt eine 10 h here maximale Taktfrequenz wenn DPR ver wendet wird 2 3 Kompiliererentwicklung 2 3 1 Frontend Das Frontend ist der erste Teil eines Kompilierungsprozesses der untersuchende Aufgaben hat Dabei wird der Quelltext
183. ure Im VIROQUANT Projekt sind die Aufgaben in drei Bereiche unterteilt A Biologie In diesem Bereich werden die Zellkulturen angelegt mit den Viren infiziert und die Zellbilder ausgewertet B Modellierung Eng in Zusammenarbeit mit Bereich A wertet dieser Bereich die Zellbilder aus und erweitert mathematische Modelle mit den Messergebnissen 1 2 Hochdurchsatzmikroskopie C Technologie Dieser Bereich arbeitet an der Verbesserung des Mikroskops und an den Bildalgorithmen zur Auswertung der Zellbilder zusammen mit den Bereichen A und B In den folgenden Abschnitten besch ftigt sich diese Dissertation weiter mit den Aufgaben aus dem Technologie Bereich C 1 2 Hochdurchsatzmikroskopie F r einen sogenannten Genomweiten Screen die Untersuchung der Signalwege bestimm ter Gene sind sehr viele Aufnahmen von Zellbildern n tig F r diese Menge ist ein Hoch durchsatzmikroskop erforderlich das sehr schnell die Zellkulturen fotografiert und mit einem Roboterarm die Proben austauscht Im Rahmen des Projekts werden existierende Mikrosko pe in ihrer Arbeitsweise beschleunigt um der Anforderung gerecht zu werden Bild 1 2 zeigt schematisch alle n tigen Komponenten f r die Datenaufnahme Sensor mit Linsen Speicher und Rechencluster N Mikroskop Roboter N FPGA Karte im PC Offline Prozessierung Bild 1 2 Hochdurchsatzmikroskop mit verbundenen Komponenten Eine Konsequenz eines schnell arbeiten
184. von Mikroskopaufnahmen mit hohem Bilddurchsatz Es wurden geeignete Klassifikations und Bildmerkmal Algorithmen studiert und ein entsprechendes Softwarepaket entwickelt Das Softwarepaket klassifiziert die Zellen in Proteinfunktionen und Ph notypen mit einer Pr zision von 83 F r die automatische Analyse der Mikroskopaufnahmen besteht ein hoher Rechenbedarf in der Gr enodnung von Monaten Das Profiling des Softwarepakets zeigt dass der Haralik Texturen Merkmal Algorithmus Ha ralick Algorithmus den gr ten Teil an Rechenzeit konsumiert Im Quelltext des Algorithmu ses wurden bereits Bem hungen unternommen die Laufzeit mit optimierenden Ma nahmen zu beschleunigen die in Kapitel 4 1 verbessert und erweitert werden Die zweite Arbeit benutzt ein FPGA auf dem die co occurrence Matrizen und die Haralick Texturen Bildmerkmale in einem speziell angepassten Hardwaredesign berechnet wurden Die Berechnungen der Matrizen deren Normierungen sowie der Bildmerkmale wurden mit Handel C in einem Design realisiert Die Ergebnisse zeigen eine Beschleunigung von 4 75 bei der Be rechnung der co occurrence Matrizen und eine Beschleunigung von 7 3 bei der Berechnung der Bildmerkmale gegen ber einer CPU Implementierung Bei dieser Entwicklung wurden lediglich sieben Bildmerkmale implementiert was f r das VIROQUANT Projekt nicht reicht denn es werden 13 Bildmerkmale f r die Auswertung der Bilder ben tigt Der verwendete FPGA hat 20 tausend Slices di
185. werden 4 Im letzten Schritt wird aus allen gesammelten und berechneten Informationen eine VHDL Datei mit der Hardware Pipeline generiert Das UML C Klassendiagramm in Bild 5 5 zeigt die Softwarestruktur des VHDL Backends Das Konzept baut auf Klassen mit Kompositionsbeziehungen auf die farbig markiert sind Jede Klasse steht in einer existenzabh ngigen Teil Von Beziehung zur dar ber liegenden d h jede Klasse h lt eine Referenz der abh ngigen Klasse die beim Konstruktoraufruf mit angegeben werden musste Aus dieser Struktur ergibt sich die Reihenfolge in der die Instanzen angelegt werden m ssen Ebenso ergibt sich die Reihenfolge der Methodenausf hrung Die Instanz von source h lt lediglich den SSA Quelltext und bernimmt die Aufgaben der le xikalischen Analyse den Quelltext in Tokens zu zerlegen Die existenzabh ngige Instanz von SourceCheck pr ft im Quelltext ob genau eine Kernelfunktion enthalten ist die Kernelfunktion keinen R ckgabewert hat daf r Parameter besitzt und ob die Syntax des Funktionsrumpfes keine Fehler aufweist Wenn die Pr fung fehl schl gt ist es garantiert dass die folgenden In stanzen ihre Aufgaben nicht ausf hren k nnen und der bersetzungsvorgang bricht an dieser Stelle ab Die n chste Instanz sourceattrib extrahiert Attribute aus dem Quelltext Ein Attri but ist die Parameterliste der Funktion die Parameterpaare listet F r jedes Paar wird der Name und der Typ eines Parameters festgest
186. windigkeitsgewinn 2 22 Cm m nn rn Skalierung des Algorithmusses 2 22 Con nn nn Optimierungsergebnisse der zweiten Version OpenCL Kompilier 2 2 2 oo nn nn 6 2 1 6 2 2 6 2 3 6 2 4 6 2 5 Nutzen der OpenCL Implementierung Nutzen des Pipelinekonzepts 2 2 2 2 2 m rn nn nn Bandbreite Speicherkontroller 2 2 o nn nn Ressourcenbedarf des Designs a 2 2 2 nn nen Beispiel Applikationen 2 2 2 on nommen 7 Fazit und Ausblick 7 1 Zieleder Arbeit a e e an au nr nn dea a a e a 7 2 Verbesserungen f r die Zukunft 2 2 2 non ee e A OpenCL FPGA Beispielanwendung Literatur XIV 111 111 111 112 113 113 113 114 115 115 116 119 119 121 123 134 1 Einf hrung und Ziele des Viroquant Projekts 1 1 Orientierung im Viroquant Projekt Weltweit sterben j hrlich drei Millionen Menschen an einer HIV Infektion Da sich 60 Millio nen Menschen im selben Zeitraum infizieren wird die Zahl an HIV Toten ohne Heilmittel stark ansteigen An chronischen Hepatitis Infektionen sterben j hrlich zwei Millionen Menschen er schreckende 500 Millionen erkranken Diese Zahlen stammen aus dem Projektantrag f r das VIROQUANT Projekt und sie alleine stellen eine ausreichende Motivation dar um aussichts reiche Forschungen zu f rdern die entweder Infizierten helfen k nnen oder Neuinfektionen reduzieren sollen Im VIROQUANT Projekt wird an den In
187. with source status status clBuildProgram program 1 amp device NULL NULL NULL if status CL_SUCCESS coutAndExit Could not build program status 123 76 77 78 79 80 81 A OpenCL FPGA Beispielanwendung kernel clCreateKernel program kernel amp status if status CL_SUCCESS coutAndExit Could not create kernel status queue clCreateCommandQueue context device 0 amp status if status CL_SUCCESS coutAndExit Could not create a command queue status f meml clCreateBuffer context CL_MEM_READ_WRITE sizeof unsigned int MEGA NULL amp status if status CL_SUCCESS coutAndExit Could not create a memory object status r mem2 clCreateBuffer context CL_MEM_READ_WRITE sizeof unsigned int MEGA NULL amp status if status CL_SUCCESS coutAndExit Could not create a memory object status r h_meml new unsigned int MEGA h_mem2 new unsigned int MEGA for i 0 i lt MEGA i h_mem1 i i h_mem2 i 0 status clEnqueueWriteBuffer queue meml true 0 sizeof unsigned int MEGA h_meml 0 NULL NULL if status CL_SUCCESS coutAndExit Could not transfer memory status status clSetKernelArg kernel 0 sizeof meml meml if status CL_SUCCESS coutAndExit Could not set kernel argument status status clSetKernelArg kernel 1 sizeof meml m
188. yps Die Bio logen k nnen Algorithmen in OpenCL schreiben und in ein Hardwaredesign f r den FPGA bersetzen was in einer Hardwarebeschreibungssprache f r sie zu komplex w re Neben der Einfachheit hat die parallele Sprache OpenCL den Vorteil der Portierbarkeit auf andere Archi tekturen Falls der FPGA Kompilierer wegen existierender Einschr nkungen den Algorithmus nicht bersetzen kann l sst sich das OpenCL Programm auch f r die GPUs in der Offlinepro zessierung bersetzen Schl sselworte CUDA Compiler Co Prozessor DPR FPGA GPGPU LLVM OpenCL Pipelinegene rator VHDL rekonfigurierbare Logik Abk rzungen ALU AST CLB CMP CPU CU CUDA DDR3 DMA DPR FIFO FPGA FPU GPGPU GPU HPC IR LLVM MEM PCIe PE SIMD SSA TCL VHDL Arithmetic Logical Unit arithmetische Recheneinheit Abstract Syntax Tree abstrakter Syntaxbaum Configurable Logic Block konfigurierbare Logikzelle Compare Vergleich Central Processing Unit Computerprozessor Compute Unit OpenCL Recheneinheit in OpenCL Compute Unified Device Architecture Programmiersprache f r Grafikkarten Double Data Rate 3 Speicherzugriffstechnik Direct Memory Access direkter Speicherzugriff Dynamic Partial Reconfiguration dynamische partielle Rekonfiguration First in First out Stapelspeicher Field Programmable Gate Array feldprogrammierbare Gatteranordnung Floating Point Unit Flie komma Recheneinheit General Purpose GPU GPU f r allgemeine Zwe

Download Pdf Manuals

image

Related Search

Related Contents

Manual del Usuario  Salon du livre final_LR - Institut Français  Utilisation du téléphone de bureau SIP Avaya E129  Reactivos - Annar Diagnóstica Import  Work Big IN35, IN35W, IN37  Quokka user manual 30 July 2012  PDF形式 ：133KB  Logik L19LID628WE User's Manual  Samsung ES25 Инструкция по использованию  iGateway user guide  

Copyright © All rights reserved.
Failed to retrieve file