OpenCl goes AutoIt Update 31.Dezember 2016

  • OpenCl goes AutoIt Update 31.Dezember 2016
    Download: OpenCL_Forum.zip
    //EDIT 31.12.2016

    • einige kleinere Bugs gefixt
    • an AutoIt Version 3.3.14.0 angepasst (Konstantenbezeichnungen)
    • DEVICES64 um Speichertransferzeiten erweitert
    • volle 64Bit-Unterstützung


    Hallo zusammen!
    Um die enorme Rechenpower aktueller CPU´s und Grafikkarten auszunutzen, war AutoIt bisher eher ungeeignet.
    Von den vielen teuer bezahlten Cores der CPU wurde maximal nur einer benutzt, die Rechenleistung der Grafikkarte kam nur einigen Grafikbibliotheken zugute.
    Manch einer kam auf die Idee, den ehrwürdigen AutoIt-Interpreter mittels Assembler- bzw. C-Bibliotheken etwas zu beschleunigen... :rock:

    OpenCl (Open Computing Language) bietet nun die Möglichkeit, die Rechenleistung von CPU, GPU und anderen Prozessoren mittels paralleler Berechnungen und portablem Code Plattformunabhängig auszunutzen!
    OpenCl ist in den neuen Treibern für Grafikkarten bereits integriert. Die SDK´s enthalten reichlich Beispielcode, weiterhin sind tolle Tools enthalten, Profiler, Kerneldebugger und vieles mehr....vollintegriert ins VisualStudio10!
    Ich verweise auf folgende Dokumente, die jeweils aktuellen Versionen erscheinen auf den Entwicklerseiten.
    OpenCL Reference Pages
    Einführung OpenCl auf deutsch
    AMD-SDK
    Nvidia-SDK
    Spezifikation
    AMD-Guide
    Die SDK´s und Treiber gibt es von allen Prozessorherstellern, aktuelle x86-CPU´s und Grafikkarten von AMD und Nvidia werden natürlich unterstützt.

    //EDIT 19.August 2012
    Intel Opencl_runtime only 2012 mehr wird für Intel-Prozessoren nicht gebraucht


    Die OpenCL-API beinhaltet reichlich Funktionen, ich habe die meisten bereits nach AutoIt portiert, aber eine UDF würde ich das im derzeitigen Zustand nicht nennen. Eher ein Proof of concept. Die 2D-Funktionen sind bereits in der Warteschleife 8o


    Was macht OpenCL aus?
    OpenCl lebt von der Parallelisierbarkeit der Algorithmen. Gut parallelisierbar sind z.B. Grafik- und Bildbearbeitungsfunktionen (Filter) oder Sortierfunktionen. Wissenschaftler verwenden bspw. Matrixmultiplikationen, Fouriertransformationen und andere gut parallelisierbare Algorithmen.
    Jedes gut in viele gleiche Teilstücke zu zerlegende programmtechnische Problem wird von OpenCL profitieren!
    Dabei verteilt OpenCL selbstständig das Programm auf alle zur Verfügung stehenden Recheneinheiten.
    Gerade die bis zu mehreren hundert Recheneinheiten auf den Prozessoren der Grafikkarten (GPU) werden ausgenutzt, was gegenüber den CPU´s teilweise immense Geschwindigkeitsvorteile bringt!
    Nicht umsonst werden in den großen Rechenzentren immer häufiger Grafikkarten zum "crunchen" eingesetzt, die CPU´s dienen größtenteils nur noch dazu, diese Grafikkartencluster zu "füttern".

    Wie kommt das Programm nun dazu, auf allen Prozessoren (Cores) gleichzeitig zu laufen? Die Threadverwaltung übernimmt glücklicherweise OpenCl, sind mehr Prozessoren verfügbar, werden diese einfach mitbenutzt...
    Das zentrale Objekt ist der sogenannte "Kernel", ein Programm in C-Syntax, welches während der Laufzeit vom Treiber kompiliert, und dann auf alle Prozessoren verteilt wird.
    Dieser Kernel ist in der Regel in einer einfachen Textdatei abgespeichert, also völlig Plattformunabhängig und somit portabel!

    Die SDK´s enthalten reichlich Beispiele (meistens in C++/C#),aber die dort verwendeten Kernel (die *.cl-Dateien) können logischerweise alle auch mit AutoIt verwendet werden!


    Los gehts mit OpenCL in AutoIt!
    Wer sich ob der Einfachheit gegenüber den Funktionen in den Beispielprogrammen wundert, ich habe viele API-Funktionsaufrufe in EINFACHEN AutoIt-Funktionen gekapselt. Wer hardcore-optimiert oder spezielle API-Funktionen benötigt, kann natürlich auch die API-Calls benutzen!

    Die folgenden Beispiele beinhalten zwar einige Erklärungen, aber wer Fragen oder Probleme hat.....wozu ist das Forum da?^^
    Problembeschreibungen bitte mit Prozessor/Grafikkarte und dem Consolenausdruck posten!
    Beispiel1, eine einfache Quadrierung von Floatzahlen

    Spoiler anzeigen
    [autoit]

    #include <opencl_inc.au3>#include <opencl_easy.au3>;OpenCl kann benutzt werden, um rechenintensive, gut parallel zu verarbeitende Programme auf Grafikkarten der neueren Generation auszuführen.;Diese Grafikkarten haben eine GPU, welche aus hunderten von einzelnen Prozessoren besteht, die wiederum unabhängig voneinander parallel arbeiten können.;Die Hardware unterteilt sich in mehrere Compute-Units, die sich wiederum in Workgroups und diese wiederum in Workitems (Shader) unterteilen.;Man kann diese vielen Prozessoren (Shader) der GPU gleichzeitig nutzen, ohne eine komplizierte Threadverwaltung (wie bei einer CPU) zu verwenden.;OpenCl verteilt das programmtechnische "Problem" selbstständig auf die verfügbaren Ressourcen.;Wenn keine von OpenCl unterstützte GPU gefunden wird, kann eine geeignete CPU die Berechnungen übernehmen. Durch Verwendung von SIMD wird auch bei einer CPU der;verwendete Code stark beschleunigt!;Um OpenCl zu nutzen, wird ein sog. "Kernel" benötigt.;Der "Kernel" ist ein Programm, dass GLEICHZEITIG auf jedem der vielen verfügbaren Shader(Prozessoren in einer GPU) der Grafikkarte läuft. Die Verwaltung der parallel;ausgeführten Threads übernimmt OpenCl.;Die im Kernel verwendeten sog. Thread-ID´s verweisen auf die Daten im Eingangs- bzw. Ausgangspuffer.;Beispiel: Es sollen 1024 Floatzahlen aus dem Eingangspuffer gelesen, im Kernel(Thread) quadriert, und anschliessend in den Ausgangspuffer geschrieben werden.;Abhängig von der Hardware werden für (fast) jede zu quadrierende Zahl ein einzelner Thread erzeugt und einem Shader zugewiesen.;Dies geschieht automatisch! Man muss sich nicht um die Verwaltung dieser Threads kümmern.;Der Kernel mit der Thread-ID=1 quadriert die 1. Zahl im Eingangspuffer und schreibt sie in den Ausgangspufferan die 1. Stelle.;Der Kernel mit der Thread-ID=17 quadriert die 17. Zahl im Eingangspuffer und schreibt sie in den Ausgangspuffer an die 17. Stelle.;Der Kernel mit der Thread-ID=1022 quadriert die 1022. Zahl im Eingangspuffer und schreibt sie in den Ausgangspuffer an die 1022. Stelle.;usw, bis alle Threads abgearbeitet sind. Die Threads werden je nach Menge der verfügbaren Shader parallel verarbeitet!;Sind weniger Shader als Daten vorhanden, werden die Shader mehrmals nacheinander benutzt, mit jeweils anderer Thread-ID.;Wird eine geeignete Grafikkarte verwendet, arbeiten dort hunderte Shader parallel und sind zig- mal schneller als eine CPU!;;Der Kernel ist eine Funktion in C-Syntax, deren Funktionsparameter im Kernel verarbeitet werden.;Der Kernel wird entweder zur Laufzeit compiliert (wie hier im Beispiel) und auf die Shader übertragen, oder als vorkompilierter Kernel verarbeitet.;Der folgende Kernel hat als Funktionsparameter den Eingangspuffer und den Ausgangspuffer und einen Parameter, welcher die Anzahl der zu berechnenden;Quadratzahlen bestimmt. Die Puffer werden über ihre Pointer (Zeiger in den Speicherbereich) angesprochen, weitere Parameter;übergibt man entweder als Zeiger oder als Wert direkt....C-Versteher haben Vorteile^^;Im Compiler ist eine Syntaxanalyse des Kernels eingebaut, bei fehlerhaftem Kernel erscheint eine Info-Box mit Beschreibung des Fehlers!Global $KernelSource = "__kernel void quadrat( __global float* inputbuffer, __global float* outputbuffer, const unsigned int count)" & @CRLF & _ "{" & @CRLF & _ "int ThreadID = get_global_id(0);" & @CRLF & _ ;hier wird von OpenCl das zu bearbeitende Element zugewiesen "if(ThreadID < count)" & @CRLF & _ ;count kann benutzt werden, um weniger Daten zu berechnen, als Shader zur Verfügung stehen "outputbuffer[ThreadID] = inputbuffer[ThreadID] * inputbuffer[ThreadID];" & @CRLF & _;holt Element aus dem Eingangspuffer, berechnet Ergebnis und schreibt in den Ausgangspuffer "}";Nun muss OpenCl bzw. die Grafikkarte auf den Kernel vorbereitet werden:;Ich habe die folgenden EINFACHEN Funktionen geschrieben um das Prinzip von OpenCl zu verdeutlichen, die Funktionen haben es in sich....;Im einfachen Beispiel hier wird die erste gefundene Grafikkarte verwendet, wer mehrere Grafikkarten (und CPU´s) parallel verwenden möchte, sollte sich;in OpenCl einarbeiten^^;;Zunächst werden die verfügbaren Geräte(Devices) gesucht._CL_GetDevice("ALL") ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.;u.a. werden auch noch die OpenCl-spezifischen Variablen Context, CommandQueue, Program ermittelt, und der Kernel kompiliert.;Wird keine geeignete Grafikkarte gefunden, wird der Kernel mit SSE-Unterstützung kompiliert und auf der CPU (auf allen Cores) ausgeführt.;Die Puffer (Buffer) werden benutzt, um die Daten an die Grafikkarte zu übergeben oder von dort zurückzubekommen.;Es gibt Eingangs- und Ausgangspuffer, welche im Hauptspeicher als Struct abgebildet werden.;Diese Speicherbereiche MÜSSEN an 16-Byte aligned Speicheradressen liegen! Da align16 bei DllstructCreate() nicht zuverlässig funktioniert,;habe ich eine eigene Funktion geschrieben, _DllStructCreate16(), welche dafür sorgt, dass die Structs mit ihren Anfangsadressen entsprechend ausgerichtet sind.;Um im Beispiel 1024 Floatzahlen zu quadrieren, benötigen wir jeweils einen Puffer:;Der Eingangspuffer übergibt die Daten an den Kernel:;dazu wird im Hauptspeicher ein Bereich mit der Anzahl der Eingangsdaten beschrieben und dieser Bereich dann zur Grafikkarte transferiert$anzahl = 1024 ;Anzahl der Floatzahlen$input_buffer = _DllStructCreate16("float[" & $anzahl & "]");Rückgabe ist "Eine Variable für den Gebrauch in DllStruct-Aufrufen", sic (aus der Hilfe^^);der Puffer wird mit $anzahl zufälliger Floatzahlen gefülltFor $i = 1 To $anzahl DllStructSetData($input_buffer, 1, Random(1, 100), $i)Next;nun wird dieser Puffer für die Bearbeitung mit OpenCl verfügbar gemacht$CL_buffer_in = _CL_CreateBuffer($input_buffer);und die Daten aus dem Hauptspeicher werden in den Speicher der Grafikkarte geschrieben_CL_WriteBuffer($CL_buffer_in, $input_buffer) ;vom Hauptspeicher in die Grafikkarte;Ausgangspuffer, hier sollen die Quadratzahlen abgespeichert werden:$output_buffer = _DllStructCreate16("float[" & $anzahl & "]");Rückgabe ist "Eine Variable für den Gebrauch in DllStruct-Aufrufen";nun wird dieser Puffer für die Bearbeitung mit OpenCl verfügbar gemacht$CL_buffer_out = _CL_CreateBuffer($output_buffer);Man könnte auch nur einen Puffer für Ein-und Ausgabe verwenden....dann überschreiben die Ergebnisse die Eingangsdaten...;Jetzt werden die Funktionsparameter (Argumente) an alle Kernelprogramme übergeben;Die Funktionsparameter werden von 0 an durchnummeriert und der Parameter übergeben;Diese Funktion verwendet nur 4Byte grosse Parameter, also dword, int, uint, float usw!!!_CL_SetArg(0, "ptr*", $CL_buffer_in) ;Kernelparameter: __global float* inputbuffer erwartet das Handle des Eingangspuffers_CL_SetArg(1, "ptr*", $CL_buffer_out) ;Kernelparameter: __global float* outputbuffer erwartet das Handle des Ausgangspuffers_CL_SetArg(2, "uint*", 100) ;Kernelparameter: const unsigned int count kann die Anzahl der zu berechnenden Zahlen begrenzen!;Jetzt lassen wir alle Kernelprogramme auf allen verfügbaren Shadern gleichzeitig laufen, OpenCl verteilt die verfügbare Anzahl Shader automatisch!;Die THREADID im Kernel ist wie ein Index in einem Array, welcher auf die einzelnen Elemente (in diesem Fall im Puffer) zeigt.;Ist keine geeignete Grafikkarte vorhanden, wird die CPU mit allen Cores zur Berechnung verwendet!;Da in diesem Fall der Kernel ein mit SSE-Unterstützung kompiliertes C-Programm ist, wird so eine fast optimale Geschwindigkeit auf der CPU erreicht!;Wird eine geeignete Grafikkarte verwendet, arbeiten dort hunderte Shader parallel und sind zig- mal schneller als die CPU!$Global_Worksize = $anzahl ;die gesamte Anzahl der zu berechnenden Workitems$Local_Worksize = 0 ;Workitems/Workgroup, legt die Auslastung der Shader fest...( davon wird noch die Rede sein^^). Wenn 0, legt der Compiler einen Wert fest_CL_RunKernel($Global_Worksize, $Local_Worksize) ;liest die Anzahl Floatzahlen aus dem Eingangspuffer und schreibt das Quadrat dieser Zahlen in den Ausgangspuffer;die berechneten Daten stehen nun im Ausgangspuffer der Grafikkarte und können in den Hauptspeicher transferiert werden_CL_ReadBuffer($CL_buffer_out, $output_buffer) ;von der Grafikkarte in den Hauptspeicher;die Daten stehen nun in der Struct im Hauptspeicher und können von dort ausgelesen und angezeigt werdenFor $i = 1 To $anzahl $in = DllStructGetData($input_buffer, 1, $i) ;eingangspuffer $out = DllStructGetData($output_buffer, 1, $i) ;ausgangspuffer ConsoleWrite("Thread-ID: " & $i - 1 & " " & $in & " " & $out & @CRLF)Next;das wars schon.....

    [/autoit]


    Die Funktion _CL_GetDevice() wird verwendet um den Kernel einem bestimmten Gerät (GPU oder CPU) zuzuweisen.
    Bei _CL_GetDevice("CPU") wird die CPU verwendet, _CL_GetDevice("GPU") benutzt den Prozessor der Grafikkarte. _CL_GetDevice("ALL") sucht die nächstbeste mögliche Hardware.
    Der extreme Unterschied zwischen Grafikkarte und CPU wird am folgenden Beispiel deutlich. Wenn eure vorhandene Grafikkarte nicht unterstützt wird, macht nix, es werden wenigstens alle Cores der CPU verwendet!
    Um in den Genuss erhöhterAusführungsgeschwindigkeit durch eine GPU zu kommen, ist eine einfache Einsteiger-Grafikkarte ausreichend. Superteuer muss sie nicht sein, ich habe mir anhand dieser Seite ein Modell in der 70€-Klasse ausgesucht, die passt zu meinem 45€-Prozessor, ist aber zwischen 10 und 100x schneller in der Berechnung einiger Kernel!
    Einige Kernel werden auch auf der CPU (wesentlich) schneller ausgeführt, wieso das so ist, wird sich im Verlauf dieses Threads noch herausstellen...

    Beispiel 2, der aus dem ASM-Tutorial bekannte Tunnelflug, von mir als Basic-Freak wohl eher dilletantisch in das Kernel-C umgesetzt.
    Man beachte den auskommentierten Teil am Ende des Scriptes.

    Spoiler anzeigen
    [autoit]

    #include <opencl_inc.au3>#include <opencl_easy.au3>#include <GUIConstantsEx.au3>#include <WinAPI.au3>#include <GDIConstants.au3>;der schon im ASM-Tut optimierte "Tunnelflug" hier in der einfachsten Version;ich bin nicht der C-Freak, bitte also um Nachsicht mit meinen Kernels, Verbesserungsvorschläge werden selbstverständlich umgesetzt^^Global $KernelSource = "__kernel void tunnelflug( __global int* output, const unsigned int width,const unsigned int height,const float stepx)" & @CRLF & _ "{" & @CRLF & _ "uint threadid = get_global_id(0);" & @CRLF & _ ;thread-id, d.h. jedes pixel von 1 bis b*h "float y = threadid/width;" & @CRLF & _ ;pixelkoordinaten "float x = threadid%width;" & @CRLF & _ ;pixelkoordinaten "float dy = (y / height) - 0.5f;" & @CRLF & _ ;...Berechnungen "float dx = (x / width) - 0.5f;" & @CRLF & _ "float wurzel = sqrt(dy * dy + dx * dx);" & @CRLF & _ "int tu = 203.718330624f * atan2(dy ,dx);" & @CRLF & _ "int tv = (((100 / wurzel) + stepx));" & @CRLF & _ "int t = (tu^tv)&255;" & @CRLF & _ ;farbe pixel "output[threadid] = 0xFF000000+(t<<16)+(t<<8)+t;" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t "}";Daten$width = 256 * 3 ;breite und höhe der gui/grafik, die Anzahl der "Threads" sollte ein Vielfaches von 16 sein!$height = 256 * 3$step = 1.02 ;das ist die Schrittweite, um die der Tunnelflug ein Frame weiter gezeichnet wird$steps = 0 ;Beginn des Fluges...Opt("GUIOnEventMode", 1)$hgui = GUICreate("", $width, $height, 1, 1) ;GUI erstellen$hdc_gui = _WinAPI_GetDC($hgui) ;HDC holen zum blittenGUISetState()GUISetOnEvent($GUI_EVENT_CLOSE, "_Exit")AdlibRegister("_fps", 1000) ;FramesPerSecond$fps = 0;bitmap erzeugen, in $ptr_bitmap steht nach dem Funktionsaufruf der Pointer auf die PixeldatenLocal $ptr_bitmap, $hbmp_bitmap ;byref$hDC_bitmap = _CreateNewBmp32($width, $height, $ptr_bitmap, $hbmp_bitmap) ;DC, Pointer auf die Bitmapdaten und ein Handle für GDI+....eine eierlegende WollmilchsauGlobal $DATA_SIZE = $width * $height ;bildgröße=anzahl pixel (rgba)Global $NULL = 0;ab hier gehts mit OpenCl los;was wird gebraucht?;Buffer für Output der Daten in der Größe der Bitmap, also;Zunächst werden die verfügbaren Geräte(Devices) gesucht._CL_GetDevice("ALL") ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.;Puffer für die Bitmap$output_buffer = DllStructCreate("dword[" & $DATA_SIZE & "]", $ptr_bitmap);halleluja! bitmaps werden immer 16byte-aligned!$CL_buffer_out = _CL_CreateBuffer($output_buffer);Parameter an den Kernel übergeben_CL_SetArg(0, "ptr*", $CL_buffer_out)_CL_SetArg(1, "uint*", $width)_CL_SetArg(2, "uint*", $height)_CL_SetArg(3, "float*", $steps)$CL_DEBUGFLAG = 0 ;Debug-Ausgabe ausschalten, um Geschwindigkeit zu erhöhen;Schleife zur Anzeige für den TunnelflugWhile 1 ;endless loop... $steps += $step ;schrittweite für den nächsten frame erhöhen... _CL_SetArg(3, "float*", $steps) ;...und an kernel übergeben _CL_RunKernel($DATA_SIZE, 0) ;Kernel ausführen _CL_ReadBuffer($CL_buffer_out, $output_buffer);Puffer(alle Pixel) lesen _WinAPI_BitBlt($hdc_gui, 0, 0, $width, $height, $hDC_bitmap, 0, 0, $srccopy) ;Bitmap in die GUI blitten $fps += 1 ;Frames pro Sekunde zählenWEnd;...das wars schon^^Func _CreateNewBmp32($iwidth, $iheight, ByRef $ptr, ByRef $hbmp) ;erstellt leere 32-bit-Bitmap; Rückgabe DC und ptr und handle auf die Bitmapdaten ;by Andy Local $hcdc = _WinAPI_CreateCompatibleDC(0) ;Desktop-Kompatiblen DeviceContext erstellen lassen Local $tBMI = DllStructCreate($tagBITMAPINFO) ;Struktur der Bitmapinfo erstellen und Daten eintragen DllStructSetData($tBMI, "Size", DllStructGetSize($tBMI) - 4) ;Structgröße abzüglich der Daten für die Palette DllStructSetData($tBMI, "Width", $iwidth) DllStructSetData($tBMI, "Height", -$iheight) ;minus =standard = bottomup DllStructSetData($tBMI, "Planes", 1) DllStructSetData($tBMI, "BitCount", 32) ;32 Bit = 4 Bytes => AABBGGRR Local $adib = DllCall('gdi32.dll', 'ptr', 'CreateDIBSection', 'hwnd', 0, 'ptr', DllStructGetPtr($tBMI), 'uint', $DIB_RGB_COLORS, 'ptr*', 0, 'ptr', 0, 'uint', 0) $hbmp = $adib[0] ;hbitmap handle auf die Bitmap, auch per GDI+ zu verwenden $ptr = $adib[4] ;pointer auf den Anfang der Bitmapdaten, vom Assembler verwendet _WinAPI_SelectObject($hcdc, $hbmp) ;objekt hbitmap in DC Return $hcdc ;DC der Bitmap zurückgebenEndFunc ;==>_CreateNewBmp32Func _fps() WinSetTitle($hgui, "", $fps & " FPS") $fps = 0EndFunc ;==>_fpsFunc _Exit() ExitEndFunc ;==>_Exit;ein Kernel, der die SIMD-Fähigkeiten vor allem der CPU´s wesentlich besser ausnutzt!;es werden 4 Pixel gleichzeitig berechnet, auf der GPU bringt das nur ca. 10-20%, aber auf der CPU ca. 50-60% mehr Geschwindigkeit!;~ Global $KernelSource = "__kernel void tunnelflug( __global int* output, const unsigned int width,const unsigned int height,const float stepx)" & @CRLF & _;~ "{" & @CRLF & _;~ "uint threadid = get_global_id(0)*4;" & @CRLF & _ ;thread-id, d.h. jedes 4. pixel von 1 bis b*h;~ "float4 y = {threadid/width,(threadid+1)/width,(threadid+2)/width,(threadid+3)/width};" & @CRLF & _ ;pixelkoordinaten;~ "float4 x = {threadid%width,(threadid+1)%width,(threadid+2)%width,(threadid+3)%width};" & @CRLF & _ ;pixelkoordinaten;~ "float4 dy = (y / height) - 0.5f;" & @CRLF & _;...Berechnungen;~ "float4 dx = (x / width) - 0.5f;" & @CRLF & _;~ "float4 wurzel = sqrt(dy * dy + dx * dx);" & @CRLF & _;~ "float4 tu = 203.718330624f * atan2(dy ,dx);" & @CRLF & _;~ "float4 tv = (100 / wurzel) + stepx;" & @CRLF & _;~ "int4 tu1 = convert_int4(tu);" & @CRLF & _;~ "int4 tv1 = convert_int4(tv);" & @CRLF & _;~ "int4 t = (tu1^tv1)&255;" & @CRLF & _ ;farbe pixel;~ "output[threadid] = (4278190080+(t.s0<<16)+(t.s0<<8)+t.s0);" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t;~ "output[threadid+1] = 4278190080+(t.s1<<16)+(t.s1<<8)+t.s1;" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t;~ "output[threadid+2] = 4278190080+(t.s2<<16)+(t.s2<<8)+t.s2;" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t;~ "output[threadid+3] = 4278190080+(t.s3<<16)+(t.s3<<8)+t.s3;" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t;~ "}";Den Kernel nach oben kopieren und;CL_RunKernel($DATA_SIZE, 0) ;Kernel ausführen;ersetzen durch;CL_RunKernel($DATA_SIZE/4, 0) ;Kernel ausführen;da nur noch jedes vierte Workitem Daten aus dem Puffer holen muss!!!

    [/autoit]

    Besteht bei (von Hand und viel KnowHow und Arbeit) hochoptimiertem Assemblercode (auf der CPU) noch ein gewisser Geschwindigkeitsvorteil gegenüber den gängigen Hochsprachen, wird auf der GPU klar, WO der Hammer hängt!
    Faktor 10-20 in der Geschwindigkeit, OHNE Optimierung! Das funktioniert allerdings nur bei gut zu parallelisierenden Algorithmen. Und Platz für Optimierungen innerhalb der Kernel ist reichlich vorhanden, bedingt durch die "andere" Struktur der Grafikprozessoren, Speicheranbindungen usw.
    Da es für alle gängigen Hochsprachen OpenCL-Unterstützung gibt, reduziert sich der Aufwand zum Beschleunigen eines Programms auf das Umschreiben des "Inner Loop" in einen Kernel, um massiv von der Geschwindigkeit von Mehrkernprozessoren bzw. Grafikkarten zu profitieren!

    Letzes Beispiel: Ein Flug durch die Mandelbrotmenge, diesen Kernel habe ich mir "ausgeliehen" (man muss das Rad nicht immer neu erfinden^^), im Vergleich dazu verweise ich auf den Assemblercode HIER.

    Spoiler anzeigen
    [autoit]

    #include <opencl_inc.au3>#include <opencl_easy.au3>#include <GUIConstantsEx.au3>#include <WinAPI.au3>#include <GDIConstants.au3>;diesen Kernel zur Berechnung der Mandelbrotmenge habe ich mir "ausgeliehen" ^^Global $KernelSource = FileRead("mandelbrot_kernels.cl")ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $KernelSource = ' & $KernelSource & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console;Daten$width = 256 * 3 ;breite und höhe der gui/grafik, die Anzahl der "Threads" sollte ein Vielfaches von 16 sein!$height = $width ; einfacher^^$cx = -1.76853670262974 ;koordinaten des Flugziels$cy = -0.000753954632767545$stepx = 4 / $width$stepy = 4 / $width$maxiter = 120 ;Start IterationenOpt("GUIOnEventMode", 1)$hgui = GUICreate("", $width, $height, 1, 1) ;GUI erstellen$hdc_gui = _WinAPI_GetDC($hgui) ;HDC holen zum blittenGUISetState()GUISetOnEvent($GUI_EVENT_CLOSE, "_Exit")AdlibRegister("_fps", 1000) ;FramesPerSecond$fps = 0;bitmap erzeugen, in $ptr_bitmap steht nach dem Funktionsaufruf der Pointer auf die PixeldatenLocal $ptr_bitmap, $hbmp_bitmap ;byref$hDC_bitmap = _CreateNewBmp32($width, $height, $ptr_bitmap, $hbmp_bitmap) ;DC, Pointer auf die Bitmapdaten und ein Handle für GDI+....eine eierlegende WollmilchsauGlobal $DATA_SIZE = $width * $height ;bildgröße=anzahl pixel (rgba)Global $NULL = 0;ab hier gehts mit OpenCl los;Zunächst werden die verfügbaren Geräte(Devices) gesucht._CL_GetDevice("ALL") ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.;Puffer für die Bitmap$output_buffer = DllStructCreate("dword[" & $DATA_SIZE & "]", $ptr_bitmap);halleluja! bitmaps werden immer 16byte-aligned!$CL_buffer_out = _CL_CreateBuffer($output_buffer);Parameter an den Kernel übergeben_CL_SetArg(0, "ptr*", $CL_buffer_out)_CL_SetArg(6, "uint*", $width)_CL_SetArg(7, "uint*", 0)$CL_DEBUGFLAG = 0 ;Debug-Ausgabe ausschalten, um Geschwindigkeit zu erhöhenWhile $maxiter < 4000 ;so lange, bis float an seine Grenzen stößt $stepx *= 0.995 ;schrittweite des frames erhöhen... $maxiter /= 0.9985 ;Iterationen erhöhen _CL_SetArg(1, "float*", $cx - ($width / 2) * $stepx) _CL_SetArg(2, "float*", $cy - ($width / 2) * $stepx) _CL_SetArg(3, "float*", $stepx) _CL_SetArg(4, "float*", $stepx) _CL_SetArg(5, "uint*", $maxiter) _CL_RunKernel($DATA_SIZE / 4, 0) ;Kernel ausführen, da im Kernel SIMD verwendet wird, nur 1/4 der Threads nötig! _CL_ReadBuffer($CL_buffer_out, $output_buffer);Puffer(alle Pixel) lesen _WinAPI_BitBlt($hdc_gui, 0, 0, $width, $height, $hDC_bitmap, 0, 0, $srccopy) ;Bitmap in die GUI blitten $fps += 1 ;Frames pro Sekunde zählenWEndMsgbox(0,"OpenCl-Demo","Ende!");...das wars schon^^Func _CreateNewBmp32($iwidth, $iheight, ByRef $ptr, ByRef $hbmp) ;erstellt leere 32-bit-Bitmap; Rückgabe DC und ptr und handle auf die Bitmapdaten ;by Andy Local $hcdc = _WinAPI_CreateCompatibleDC(0) ;Desktop-Kompatiblen DeviceContext erstellen lassen Local $tBMI = DllStructCreate($tagBITMAPINFO) ;Struktur der Bitmapinfo erstellen und Daten eintragen DllStructSetData($tBMI, "Size", DllStructGetSize($tBMI) - 4) ;Structgröße abzüglich der Daten für die Palette DllStructSetData($tBMI, "Width", $iwidth) DllStructSetData($tBMI, "Height", -$iheight) ;minus =standard = bottomup DllStructSetData($tBMI, "Planes", 1) DllStructSetData($tBMI, "BitCount", 32) ;32 Bit = 4 Bytes => AABBGGRR Local $adib = DllCall('gdi32.dll', 'ptr', 'CreateDIBSection', 'hwnd', 0, 'ptr', DllStructGetPtr($tBMI), 'uint', $DIB_RGB_COLORS, 'ptr*', 0, 'ptr', 0, 'uint', 0) $hbmp = $adib[0] ;hbitmap handle auf die Bitmap, auch per GDI+ zu verwenden $ptr = $adib[4] ;pointer auf den Anfang der Bitmapdaten, vom Assembler verwendet _WinAPI_SelectObject($hcdc, $hbmp) ;objekt hbitmap in DC Return $hcdc ;DC der Bitmap zurückgebenEndFunc ;==>_CreateNewBmp32Func _fps() WinSetTitle($hgui, "", $fps & " FPS") $fps = 0EndFunc ;==>_fpsFunc _Exit() ExitEndFunc ;==>_Exit

    [/autoit]


    Eins möchte ich noch betonen, die Beispiele sind auf "Einfachheit" getrimmt, es ist natürlich vollkommen blödsinnig, eine Grafik mittels OpenCl auf einer Grafikkarte "rendern" zu lassen, um dann diese Grafik in den Hauptspeicher zu transferieren (DAS kostet ca die 100- bis 1000fache Zeit wie das Rendern an sich!!!) und von dort per GDI wieder zur Darstellung an die Grafikkarte zu schicken....
    OpenGL ist das Stichwort, die üblichen Verdächtigen hier im Forum werden sicherlich passende Beiträge einbringen :thumbup:


    Ich hoffe auf Feedback und einige interessante Vorschläge für Kernel! Viel Spass mit OpenCl!

    P.S.: Großen Dank an TheShadowAE, Sprenger120 und AspirinJunkie fürs Testen auf den diversen Hardware-Plattformen!!!

    //EDIT 8. November 2011:
    Kleine Bugfixes, Log enthält bei Inteltreibern auch ohne Fehler einen Inhalt, Danke an UEZ fürs Testen

    //EDIT 19.August 2012
    Die einige Files sind teilweise angepasst/modifiziert/weiterentwickelt.
    ab diesem Post befinden sich die Vergleiche CPU/GPU als Ergebnis der Mandelbrot-Berechnung

    //EDIT 26. August 2012
    Sollte jetzt auch auf 64Bit-Maschinen laufen
    OpenCl64.au3 ermittelt die vorhandenen Geräte (Devices), mit denen OpenCL-Kernel berechnet werden können.

    //EDIT 31.Dezember 2016
    Funktionen weiter angepasst


    Damit die Bestimmung des Geräts und die Kompilierung des Kernels nicht so aufwendig sind, habe ich diese in eine Funktion gewrappert.
    Mit der Funktion _CL_GetDevice() kann man die Wahl des Devices entweder selbst bestimmen per

    [autoit]

    _CL_GetDevice()

    [/autoit]

    , dann wird das nächste verfügbare Device verwendet

    [autoit]

    _CL_GetDevice("CPU")

    [/autoit]

    verwendet die CPU, GPU die GPU....
    In der Console erscheinen in blauer Schrift die verfügbaren Geräte in einem String, z.B.

    Zitat

    >Device verfügbar= 1;2;2;CPU;AMD Athlon(tm) II X2 250 Processor;129377568;89564448


    die erste Ziffer bezeichnet die Platform (INTEL/AMD/NVIDIA o.a.), die zweite Zahl das dort enthaltene Gerät (Device)
    Wenn man genau dieses Gerät verwenden möchte, dann sollte man

    [autoit]

    _CL_GetDevice("ALL",Ziffer_Device,Ziffer_Platform) ;angeben;im vorliegenden Beispiel also_CL_GetDevice("ALL",2,1);verwendet das 2. Device von der 1. Platform

    [/autoit]


    die 3. Ziffer des Strings vom verfügbaren Device bezeichnet den Typ (GPU/CPU), dann folgt der Name des Geräts, die Device_ID und die Platform_ID

    Nach dem Aufruf von _CL_GetDevice() muss man nur noch die Kernel-Parameter setzen und den Kernel laufen lassen, s. weitere Beispiele.

    OpenCL_Forum.zip

    ciao
    Andy


    "Schlechtes Benehmen halten die Leute doch nur deswegen für eine Art Vorrecht, weil keiner ihnen aufs Maul haut." Klaus Kinski
    "Hint: Write comments after each line. So you can (better) see what your program does and what it not does. And we can see what you're thinking what your program does and we can point to the missunderstandings." A-Jay

    Wie man Fragen richtig stellt... Tutorial: Wie man Script-Fehler findet und beseitigt...X-Y-Problem

    21 Mal editiert, zuletzt von Andy (6. Januar 2014 um 01:32)

  • HIER
    alternativ:

    [autoit]

    Global Const $DIB_RGB_COLORS = 0
    Global Const $SRCCOPY = 0x00CC0020

    [/autoit]
  • Zitat

    Edit: Bei $num_platforms = 0 solltest du noch einen Fehler werfen, statt deine UDF zu crashen


    Klugsch**** ;)
    Wenn ich nur 5% von deinem Skill hätte, dann WÄRE das jetzt eine UDF, und nicht ein armseliger Haufen halbgarer Funktionen :thumbup:
    Aber ich hab da was für dich..
    autoit.de/wcf/attachment/14534/
    ...die müssten alle noch angepasst und getestet werden^^

  • Zitat von progandy

    Super, Du solltest aber noch doe globale Variable $KernelSource durch einen Funktionsparameter ersetzen


    wieso, der zieht sich den Funktionsnamen doch aus dem Kernel direkt

    [autoit]

    ;Kernel erstellen für jedes einzelne workitem
    ;kernelname extrahieren, der Kernelname steht idR vor der öffnenden Klammer, Regex needed^^
    $kname1 = StringInStr($KernelSource, "__kernel")
    $kname2 = StringInStr($KernelSource, "(", 0, 1, $kname1)
    $kname = StringSplit(stringstripws(StringMid($KernelSource, $kname1, $kname2 - $kname1),6), " ", 3)
    $kernel = clCreateKernel($program, $kname[UBound($kname) - 1], $errcode_ret)

    [/autoit]
  • Hi Andy,

    Wieder mal eine saugeile Idee :thumbup:
    Nur ein Problem tut sich auf:

    Bei mir wird angeblich kein Device (weder GPU noch CPU) gefunden.

    GPU: NVIDIA 9800GTX+ (BlackMod) (1 GByte)
    CPU: Intel Core 2 Duo, 2 x 1.75 GHz (ErsatzCPU, meine normale ist K.O gegangen)

    In Bezug auf Fehlercodes kannst du Sprenger fragen, weil Ich hab die Code's jetzt nicht zur Hand.

    mfg,
    Dominik

  • Hi,

    Zitat von Greenhorn

    Warum tut Ihr beiden Andys euch nicht zusammen und optimiert das zu einer UDF ?!

    Na, schaumamal^^.
    Ich fand ja die Idee sehr interessant, mit einer Scriptsprache "ernsthaft" die Möglichkeiten aktueller Hardware voll ausnutzen zu können.
    Ausserdem hast du mich ja "heiss" gemacht, wie soll man sonst einen Bulldozer mit 8 Kernen mittels AutoIt füttern ohne sich die Finger mit Threadverwaltung usw zu brechen?
    Weiterhin gäbe es nun auch mal die Möglichkeit, gewisse UDF´s wie z.B. die BigNum etwas "aufzumöbeln"....und nicht zuletzt wurde ich gezwungen, mich mit C auseinanderzusetzen (mit dem ++-Gedöns können sie mir (vorerst) aber trotzdem den Buckel runterrutschen :rofl: )

    Blume, bitte Info, ob und wie du Opencl zum Laufen bekommen hast!
    Die neuesten Grafikkartentreiber bringen OpenCl-Unterstützung mit, ansonsten sollte zumindest der Prozessor unterstützt werden (s. Progandys Link)

  • Hi,
    Ausserdem hast du mich ja "heiss" gemacht, wie soll man sonst einen Bulldozer mit 8 Kernen mittels AutoIt füttern ohne sich die Finger mit Threadverwaltung usw zu brechen?


    :whistling:

    Funzt OpenCL auch mit nVidia-Karten ? *dummfrag*


    Gruß


  • Leider funzt es auf meinem Schlepptop mit einer Intel Gfx Karte nicht :(

    Gruß,
    UEZ

    Auch am Arsch geht ein Weg vorbei...

    ¯\_(ツ)_/¯

  • Hi,

    Zitat von Greenhorn

    Funzt OpenCL auch mit nVidia-Karten ?

    ja, bei Verwendung der neuesten Treiber (R280.xx ??) ist OpenCl-unterstüzung implementiert. Ob deine Grafikkarte aber mitspielt, weiss ich nicht^^

    Zitat von UEZ

    Leider funzt es auf meinem Schlepptop mit einer Intel Gfx Karte nicht

    Leider funzt es auf meinem PIII-Celeron mit 1,2Ghz auch nicht^^ Aber auf dem kann man noch sehr schnelle Assemblerprogramme schreiben :rofl:
    Aber Spass beiseite, ggf. funktioniert ja die CPU als Device.

  • Ich bekomme immer Fehlermeldungen (Returncode 1). Z. B. bei der einfachen Quadrierung:

    Code
    C:\Coding\AU3\GFX\OpenCI\opencl_easy.au3 (38) : ==> Array variable has incorrect number of subscripts or subscript dimension range exceeded.:
    DllStructSetData($context_properties, 1, $platform[1], 2)
    DllStructSetData($context_properties, 1, ^ ERROR

    Gruß,
    UEZ

    Auch am Arsch geht ein Weg vorbei...

    ¯\_(ツ)_/¯

  • Die eigentliche Abfrage eines Fehlers geschieht mittels der Auswertung des Error-Returncodes des API-Calls.
    Wenn du in der Msgbox die Info "DllCall failed" bekommst, dann ist der Rückgabewert der vom DllCall(), also bei 1 => "@error: 1 Die Dll-Datei kann nicht verwendet werden, "
    ggf 32-Bit-Problem? OpenCl benutzt nur 32Bit-Dll´s!

  • UEZ: Der Fehler kommt solange kein OpenCL-Gerät gefunden wird. Dann musst du je nach CPU das passende SDK installieren und neustarten. Vielleicht hilft das auch bei ein paar Grafikkarten, aber ich beweifle das. Wenn die Graka nicht will, dann hilft auch das SDK nichts. Bei der CPU sieht das aber anders aus ;)