[OpenCL] Einfaches Rechnen auf GPU

    Hey Leute,

    Da auf Win8 Developer endlich mein OpenCL klappt, hab Ich was feines gebastelt.
    Dieses Script rechnet die Grund-Rechenarten (Plus, Minus, Geteilt, Mal) auf der GPU.

    Das Script ist extrem einfach gehalten, Ich konnte es besser auch nicht :D

    Kernel
    Code
    __kernel void Rechnen(__global float* iInputInteger1, __global float* iInputInteger2, __global float* iOutputBuffer) {
	int iThreadID = get_global_id(0);
	iOutputBuffer[iThreadID] = iInputInteger1[iThreadID] [RECHEN_MODUS] iInputInteger2[iThreadID];
}
    AutoIt Quellcode
    [autoit]


    #include ".\Includes\OpenCL.au3"

    [/autoit] [autoit][/autoit] [autoit]

    $RechenArt = "+" ;Rechenart: + - / *

    [/autoit] [autoit][/autoit] [autoit]

    $KernelSource = FileRead(@ScriptDir & "\Kernel.cl") ;Kernel Code lesen
    $KernelSource = StringReplace($KernelSource, "[RECHEN_MODUS]", $RechenArt) ;Rechenart ersetzen, erspart für jeden Modus eine Datei

    [/autoit] [autoit][/autoit] [autoit]

    _CL_GetDevice("GPU") ;GPU auswählen für CL

    [/autoit] [autoit][/autoit] [autoit]

    $iRuns = 100 ;Durchläufe

    [/autoit] [autoit][/autoit] [autoit]

    $iInputBuffer1 = _DllStructCreate16("float[" & $iRuns & "]") ;Eingabe-Buffer
    $iInputBuffer2 = _DllStructCreate16("float[" & $iRuns & "]") ;Eingabe-Buffer

    [/autoit] [autoit][/autoit] [autoit]

    For $iX = 1 To $iRuns ;Beispiel-Werte setzen (1 bis $iRuns) Kann ja geändert werden
    DllStructSetData($iInputBuffer1, 1, $iX, $iX)
    DllStructSetData($iInputBuffer2, 1, $iX, $iX)
    Next

    [/autoit] [autoit][/autoit] [autoit]

    $hCL_Buffer_In1 = _CL_CreateBuffer($iInputBuffer1) ;CL Buffer erstellen
    $hCL_Buffer_In2 = _CL_CreateBuffer($iInputBuffer2) ;CL Buffer erstellen

    [/autoit] [autoit][/autoit] [autoit]

    _CL_WriteBuffer($hCL_Buffer_In1, $iInputBuffer1) ;Daten von DllStruct in CL Buffer schreiben
    _CL_WriteBuffer($hCL_Buffer_In2, $iInputBuffer2) ;Daten von DllStruct in CL Buffer schreiben

    [/autoit] [autoit][/autoit] [autoit]

    $iOutputBuffer = _DllStructCreate16("float[" & $iRuns & "]") ;DllStruct (Ausgabe) erstellen
    $hCL_Buffer_Out = _CL_CreateBuffer($iOutputBuffer) ;CL Buffer erstellen mit Daten von Struct ^

    [/autoit] [autoit][/autoit] [autoit]

    _CL_SetArg(0, "ptr*", $hCL_Buffer_In1) ;Buffer 1 an Kernel übergeben
    _CL_SetArg(1, "ptr*", $hCL_Buffer_In2) ;Buffer 2 an Kernel übergeben
    _CL_SetArg(2, "ptr*", $hCL_Buffer_Out) ;Buffer Ausgabe an Kernel übergeben

    [/autoit] [autoit][/autoit] [autoit]

    _CL_RunKernel($iRuns, 0) ;Kernel starten
    _CL_ReadBuffer($hCL_Buffer_Out, $iOutputBuffer) ;Ausgabe-Buffer lesen

    [/autoit] [autoit][/autoit] [autoit]

    For $iX = 1 To $iRuns
    $iIn1 = DllStructGetData($iInputBuffer1, 1, $iX) ;Zahl 1
    $iIn2 = DllStructGetData($iInputBuffer2, 1, $iX) ;Zahl 2
    $iOut = DllStructGetData($iOutputBuffer, 1, $iX) ;Ergebnis

    [/autoit] [autoit][/autoit] [autoit]

    ConsoleWrite(StringFormat("Ergebnis: %i " & $RechenArt & " %i = %i", $iIn1, $iIn2, $iOut) & @LF) ;Ergebnis schreiben
    Next

    [/autoit]

    Und das Paket als Download: autoit.de/wcf/attachment/14545/

    Ich hoffe euch gefällt es :)

    mfg,
    Dominik

    Hi!
    :thumbup:
    Genau SO fängt es an, erst spielt man bissl damit (Opencl) rum um mal ein Gefühl für die Sache zu bekommen, und später führt die Grafikkarte schnelle Funktionen aus...

    Wer hat Ideen, welche bestehende (langsame) AutoIt-Funktion bzw. UDF eine "Optimierung" vertragen könnte?

    Schöner Anfang, mach so weiter und du wirst noch was nützliches rausbekommen. Ich hätte ein Fernziel für dich: Nehme eine Gleichung im Stringformat an, teile sie in parallel verarbeitbare Teile und führe das ganze per OpenCL aus ;)

    Andy: _ArraySort oder _ArraySearch wäre ein perfekter Kandidat, wenn man keine DLLStructs erstellen müsste (BitonicSort)

    Zitat von Andy

    Wer hat Ideen, welche bestehende (langsame) AutoIt-Funktion bzw. UDF eine "Optimierung" vertragen könnte?

    - Funktion ähnlich wie _WinAPI_AlphaBlend aber mit (bi)kubischer Interpolation.
    - Blur (Gaussche Unschärfe)

    Ich denke diese Funktionen würden stark profitieren, da im Prinzip eine Schleife durchläuft und alle Pixel abzählt (in beiden Fällen) und bearbeitet. Das lässt sich bestimmt gut paralellisieren.
    Und man könnte (sofern der Alphablendersatz schnell ist) vllt auch mal größere/mehr Grafiken nutzen ohne direkt auf 100% Programmauslastung zu kommen.

    Kommt man irgendwie an den Grafikpuffer für das Ausgabebild welches auf den Schirm gelangt ?
    Dann könnte man vllt sowas wie eine Bitmap im Grafikram anlegen die dann extrem schnell bearbeitet und angezeigt/skaliert usw. werden kann.
    Am besten wären alle Verwendeten Grafiken in einem Spiel komplett im Grafikram. Dann müsste nicht mehr so viel über den Bus.

    lg
    M

    Zitat von progandy

    Andy: _ArraySort oder _ArraySearch wäre ein perfekter Kandidat, wenn man keine DLLStructs erstellen müsste (BitonicSort)


    ja, aber das müsste man einfach mal ausprobieren. Letztendlich ist einfach nur die Frage, ab welcher Listengröße sich der overhead von der Konvertierung des Arrays in eine Struct, die anschließende Sortierung und das Zurückschreiben in ein Array lohnt...
    Um Arrays ganz zu vermeiden, hatte ich aus genau dafür mal eine Funktion geschrieben, die bestehenden "Array"-Code in "DllStruct"-Code umwandelt, d.h. den kompletten AutoIt-Quellcode von Array nach Struct umschreibt. Das war eigentlich recht simpel und hatte gut funktioniert. Ich wollte einfach mal die Geschwindigkeit vergleichen, und die war in einem "normalen" AutoIt-Programm in etwa gleich, daher habe ich mich nicht weiter hineingekniet.
    Mit Hinblick auf OpenCl sieht das natürlich schon anders aus....
    Gerade bei Programmen, welche große Arrays oft sortieren bzw. im Array suchen müssen, wäre der Geschwindigkeitsvorteil eklatant!

    Zitat von marsi

    - Funktion ähnlich wie _WinAPI_AlphaBlend aber mit (bi)kubischer Interpolation.
    - Blur (Gaussche Unschärfe)
    Ich denke diese Funktionen würden stark profitieren, da im Prinzip eine Schleife durchläuft und alle Pixel abzählt (in beiden Fällen) und bearbeitet. Das lässt sich bestimmt gut paralellisieren.
    Und man könnte (sofern der Alphablendersatz schnell ist) vllt auch mal größere/mehr Grafiken nutzen ohne direkt auf 100% Programmauslastung zu kommen.

    Sorry marsi, hatte deinen Post glatt überlesen...ja, sämtliche gängigen Grafikfilter lassen sich durch die Parallelisierung extrem beschleunigen.
    Mit dem netten Nebeneffekt, daß die CPU sich in den "Schlafmodus" begibt, und daraufhin das Programm langsamer wird! Kein Witz, wenn ich auf meinem AMD-Prozessor QnQ aktiviere und per AMD Power Monitor die Aktivität der Kerne überwache stelle ich z.B. beim Tunnelflug fest, daß der Prozessor (CPU) sich von 3200 auf 800 Mhz runtertaktet und somit die FPS um ca. 20-30% einbrechen...
    Die CPU schiebt die Daten zur Graka, legt sich schlafen, wird per Aktivität auf dem PCI-Bus durchs zurückschreiben ins RAM wieder geweckt, usw usf.
    Deaktiviere ich QnQ läuft die CPU permanent mit "Vollgas" und die FPS stellen das obere Limit dar.
    Das ist natürlich nur, weil ich für das Beispiel das blödsinnige blitten verwende, d.h. die Graka schiebt die Ergebnisse der Berechnung ins RAM und die CPU schiebt per "Blit" die Grafik wieder zur Graka....
    Ergo kommen wir zu deiner nächsten Frage^^

    Zitat

    Kommt man irgendwie an den Grafikpuffer für das Ausgabebild welches auf den Schirm gelangt ?
    Dann könnte man vllt sowas wie eine Bitmap im Grafikram anlegen die dann extrem schnell bearbeitet und angezeigt/skaliert usw. werden kann.
    Am besten wären alle Verwendeten Grafiken in einem Spiel komplett im Grafikram.

    Yepp, SEHR einfach über OpenGl.
    Erstellst du die Grafik bzw. das Fenster in OpenGl, kannst du diese Daten (Grafiken) entweder direkt über die Shader per OpenGl oder aber ohne weiteres direkt mit OpenCl bearbeiten.

    Zitat

    Dann müsste nicht mehr so viel über den Bus.

    hehe, dann geht NICHTS mehr über den Bus bis auf die Handvoll Steuerbytes vom Treiber....also kein Blitten mehr, kein GDI-Gedöns und natürlich auch schnarchlangsame Grafik, wenn z.B. auf Netbooks die Grafik"hardware" emuliert werden muss.....aber einen Tod muss man immer sterben^^

    Ich würde gerne das passende Beispiel OpenGL & OpenCl in AutoIt vorstellen, der Tunnelflug sollte dann 1000 FPS locker erreichen. Die reinen Ausführungszeiten des Tunnelflug-Kernels auf der Graka sind 0.0irgendwas Millisekunden. Die Schieberei der Daten über den Bus ins RAM und von dort wieder zur Graka (blitten) dauert 100x bis 1000x so lange!!!
    Leider kann ich zzt. nur extrem wenig in Sachen AutoIt machen, der Tag hat nun mal nur 28h :D
    Aber ich bleibe am Ball 8o (ausserdem hat mein Datenkompressions-Kernel erste Priorität, da kombiniere ich einige der gängigen Verfahren mit eigenen Ideen)

    Nabend!
    Yeah, ich liebe es die Posts von Andy zu lesen - sind immer super informativ und hoch Technisch :D So wie ich das rausgelesen habe könnte man viel von dem WinAPI und GDI+ Krams doch tatsächlich beschleunigen... Ich bin mal gespannt wie sich das mit OpenCL weiter entwickelt!

    Grüsse von
    Techmix!

    Und schon sind umfangreichere Spiele die ordentliche Grafiken nutzen, sowie verschiedene Effekte in Echtzeit können (Unschärfe, Bewegungsunschärfe, Helligkeit/Kontrast usw) keine Utopie mehr, sondern evtl in absehbarer Zeit realität. Ich denke der GDI+/API Kram wird nicht beschleunigt, sondern ersetzt.

    Wahrscheinlich müsste man eine komplett neue UDF herstellen, die die GDI+ Befehle (oder äquivalente) beinhaltet, diese aber komplett per OpenCL/GL ausführt.

    Sollte jemand auf die Idee kommen dies zu tun bin ich gerne bereit zu helfen, sofern dazu kein großes Wissen über OpenCL benötigt wird. Funktionen wie z.B. ein Rechteck füllen oä traue ich mir schon zu (nachdem ich mir OpenCL mal angesehen habe :P ). Aber einen Grafikpuffer (analog zu GDI+Graphicscreatefromhwnd) oder Ähnliches werde ich nicht hinbekommen...

    Ich freue mich aber schon auf die Zukunft. Mal sehen zu was das langsame AutoIt noch so befähigt wird :P

    lg
    M

    Zitat von Techmix

    So wie ich das rausgelesen habe könnte man viel von dem WinAPI und GDI+ Krams doch tatsächlich beschleunigen... Ich bin mal gespannt wie sich das mit OpenCL weiter entwickelt!

    Beschleunigen nur im Sinne von ersetzen!

    Zitat von Marsi

    Wahrscheinlich müsste man eine komplett neue UDF herstellen, die die GDI+ Befehle (oder äquivalente) beinhaltet, diese aber komplett per OpenCL/GL ausführt.

    Genau so siehts aus....
    WinAPI / GDI "emuliert" ja auf unterstem Level die BIOS-"Grafik"-Funktionen. Jeder noch so simple Kram wird komplett von der CPU ausgeführt.
    Man müsste es mal testen, aber ich bin überzeugt, ein XP würde auf einem 80386 (nicht sx, der konnte nur 24 Bit^^) mit integrierter VGA (640x480x16) starten und laufen...
    Sämtliche "Grafik" (d.h. alles ausser dem Textmode) wird auf der CPU berechnet und als Bytes in den Speicher (RAM) geschrieben, drei Digital/Analogwandler (DA) wandeln (Converieren) die Bytes im Speicher zu Rot-Grün-Blau-Monitorsignalen. Den dafür zuständigen Chip nennt man daher RAMDAC (für analoge Monitore)

    Zitat von Marsi

    Funktionen wie z.B. ein Rechteck füllen oä traue ich mir schon zu

    hehe, ich traue dir da noch wesentlich mehr zu^^, aber das Beispiel ist gut!
    Ein Rechteck ist ein Haufen Pixel, nehmen wir mal der Einfachheit halber an, 1000x1000 Pixel sollen ROT dargestellt werden (0x000000FF weil ABGR).

    - Vorgehen "Rechteck füllen" per GDI (seit Windows 1.0):
    Die CPU arbeitet eine Schleife ab, startet am Anfang des für die Grafik reservierten Speichers und rast durch das RAM. Mindestens eine Million (1000x1000) mal nacheinander wird geprüft, ob sich das "Pixel", also die zzt. aktuelle Speicherstelle im RAM, im Rechteck befindet, wenn ja, dann wird 0x000000FF an die Speicherstelle geschrieben, ansonsten bleibt die Speicherstelle unverändert. Ist der Speicher beschrieben, weist das GDI wiederum über die CPU den Treiber der Grafikkarte an, bitteschön die komplette Grafik über den PCI-Bus ins RAM der Grafikkarte zu "kopieren", damit dort der RAMDAC....na das wisst ihr ja jetzt schon^^
    Aufwand für die CPU, niedrig geschätzte 3-5 Millionen Takte....pro Frame, bei "flimmerfreien" 30FPS muss die CPU pro Sekunde gewaltig ackern...

    - Vorgehen "Rechteck füllen" per OpenGl, DirectX (D3D) oder anderen hardwarebeschleunigten Funktionen:
    Die CPU sendet an den Treiber der Grafikkarte den Befehl: "Fülle Rechteck von x0,y0 bis x1,y1 mit 0x000000FF". Geschätzter Aufwand für die CPU, 2000 Takte (1000fach schneller!)
    Die moderne Grafikkarte hat nicht nur ein winziges eigenes RAM und einen RAMDAC, sondern eine gewaltig großes RAM und mehrere hundert Prozessoren neben Recheneinheiten (in Hardware), welche ausschliesslich für u.a. den Befehl "Fülle Rechteck von x0,y0 bis x1,y1" oder "Zeichne Linie von x0,y0 bis x1,y1 " oder "Fülle Kreis mit Mittelpunkt x0,y0 und Radius r" gebaut wurden! Da in Hardware realisiert, dauert die Ausführung solcher Befehle nur eine Handvoll Takte. Dank des großen Speichers wird auch nichts mehr kopiert, sondern einfach der zzt. darzustellene Bildschirmbereich im RAM gemapped, also die Startadresse innerhalb des Grafikrams geändert.
    Für spezielle, nicht in Hardware gegossene Funktionen bleiben die vielen parallel arbeitenden, programmierbaren Prozessoren, die sog. Shader. Es wird unterschieden in Pixelshader um Bildpunkte in der Farbe zu ändern, Vertexshader um komplexe Objekte zu berechnen bzw in der Grafik zu verändern und neuerdings auch Geometrieshader, um Objekte und Geometrien der Grafik hinzuzufügen. Alle diese Shader können parallel arbeiten, um also beispielsweise 50 Dreiecke auf den Bildschirm zu zeichnen wären genauso viele Takte nötig wie für 2 Dreiecke.

    OpenCl benutzt weitere Recheneinheiten der GPU, die Processing Units.
    Z.B. komplexe Filter werden blitzschnell berechnet, da die Grafikdaten nicht mehr aus bzw in den Hauptspeicher kopiert werden müssen, sondern das RAM der Grafikkarte nicht mehr verlassen. Sämtliche Berechnungen laufen komplett auf den Prozessoren der Grafikkarte, für rechen- und speicherintensive Kernel steht jedem Shader bzw Processing Unit eigenes lokales (ultraschnelles) RAM zur Verfügung. Die CPU hat somit weniger zu tun und es bleiben Kapazitäten frei für bspw. aufwendige Berechnungen, die sich schlecht parallelisieren lassen.

    Zitat

    Aber einen Grafikpuffer (analog zu GDI+Graphicscreatefromhwnd) oder Ähnliches werde ich nicht hinbekommen...

    brauchst du auch nicht, denn OpenGl stellt diese Funktionen alle schon zur Verfügung.

    Zitat

    Mal sehen zu was das langsame AutoIt noch so befähigt wird

    hehe, AutoIt bleibt das einfache Interface für schnelle und kurze Scripte. Man hat nun die Möglichkeit, Funktionen bzw Berechnungen extrem zu beschleunigen, ohne auf hochoptimierten C++- oder Assemblercode zurückgreifen zu müssen.
    Der Haken dabei ist, nur selten wird man ohne "knowhow" diese Geschwindigkeit erreichen. Beispiel Arraysort(). Da definitiv kein Zugriff über Pointer auf die Datem im "Array" möglich ist, wird man zwangsläufig sein Script auf DllStructs aufbauen müssen um OpenCl zu nutzen. Das macht den Code unleserlich und "schwerfällig". Der Vorteil von AutoIt ist dahin....
    Ob man nun unleserlichen AutoItcode debuggen muss oder unleserlichen C++-Code, ist egal^^, aber die gängigen Compiler haben schon sauschnelle hochoptimierte CPU-Bibliotheken, damit wird der Abstand zu GPU-OpenCl wesentlich geringer!