OpenCl goes AutoIt Update 31.Dezember 2016

  • Nochmal mein System detalliert:

    Standart:
    MB Asus P6T So.1366, P6T-SE mit CrossFlash Bios zu P6T - wegen SLI-Nutzung, 2x PciE x16 (Max 36 Lanes auf 3x Pci x16)
    CPU i7-920 Bloomfield, 4 Kerne 8 Threads, 4x256Kb-L2 8MB-L3 (2660Mhz, 133x20, 2793Mhz Turbo)
    GPU 2x Zotac GTX 460 AMP! 1GB-DDR5 (810Mhz GPU, 1620Mhz Shader, 2000Mhz Ram)
    RAM 6GB DDR3-1066 Kingston DualChanel RAM 3x2GB (7/7/7-20)
    NT 850W BeQuiet Netzteil
    Und eben noch 2 HDD´s...
    [Blockierte Grafik: http://i61.tinypic.com/2db9xdi.png]

    Overclocked:
    CPU @3200Mhz, 160*20, 3360Mhz Turbo
    RAM @1282Mhz (8/8/8-24)
    GPU @870/1740/2200
    CPU Overclocked:
    [Blockierte Grafik: http://i58.tinypic.com/20acf1u.png]
    CPU & GPU Overclocked:
    [Blockierte Grafik: http://i62.tinypic.com/jsja69.png]

    Andy: Ah, danke für die Erklärung! ...langsamm begreife ich was der gravierende Unterschied zwischen Threading und OpenCL ist.
    Das die Unterschiede der Architekuren bei CPU und GPU Auswirkungen auf bestimmte Berechnungen nimmt - Dachte ich mir bereits.
    Deswegen hatte ich auch gefragt ob CPU & GPU in OCL auch gemeinsam Arbeiten können (ALL), aber ist es auch möglich mehrere OCL-Kernel gleichzeitig laufen zu lassen? Also ob einer auf GPU und ein weiterer auf CPU zeitgleich Starten und unterschiedliche Aufgaben Abarbeiten können - oder begrenzt das der interne DLL-Call 'opencl.dll'?
    Ohne das ich jetzt auch nur einen Blick auf die SDK´s geworfen habe, entstehen schon viele kleine Ideen für nette Funktionen!
    Werde mich mal am WE mehr in die Materie einlesen, und dann einige Versuche mit OpenCL starten.

    Aber eins muß ich noch Fragen, ich weiß das der OCL-Kernel zur Laufzeit Kompiliert wird, wird dieser jetzt Zentral 'verwaltet' oder bekommt jede 'Recheneinheit' seinen eigenen Kernel zugewiesen - oder sogar individuell kompilierte Kernel???

    Grüsse!

    2 Mal editiert, zuletzt von Techmix (12. März 2014 um 12:24)

  • Hi,
    jaja, mach uns nur die Nase lang mit deinem Equipment :D

    Zitat

    Aber eins muß ich noch Fragen, ich weiß das der OCL-Kernel zur Laufzeit Kompiliert wird, wird dieser jetzt Zentral 'verwaltet' oder bekommt jede 'Recheneinheit' seinen eigenen Kernel zugewiesen - oder sogar individuell kompilierte Kernel???

    Der genau gleiche Kernel wird von ALLEN an der Berechnung beteiligten Workunits GLEICHZEITIG ausgeführt!
    Daher muss man, wenn man fortgeschritten genug ist ^^, zusehen, dass möglichst viele Workunits "load" haben!

    Zitat

    Deswegen hatte ich auch gefragt ob CPU & GPU in OCL auch gemeinsam Arbeiten können (ALL), aber ist es auch möglich mehrere OCL-Kernel gleichzeitig laufen zu lassen?

    ja, man kann auf jedem der Devices unabhängig voneinander gleichzeitig verschiedene Kernel laufen lassen, oder aber auch OpenCL anweisen, einen Kernel auf alle Devices zu verteilen.

    In den Beispieldateien vom SDK sind viele Anwendungen, ich habe aber versucht, mit einer handvoll gewrapperten AutoIt-Funktionen auszukommen. Ehrlich gesagt frage ich mich was einen C(++)-Programmierer dazu bringt, sein Hirn soweit abzuschalten, jedes Mal hunderte Zeilen fast identischen Code reinzukloppen. Aber scheinbar werden die Programmierer für die SDK-Beispiele nach Zeilen bezahlt^^

    Da du Interesse hast und AutoIt sich imho sehr gut eignet, OpenCL-Funktionen zu wrappern, würde ich mich jedenfalls freuen, wenn du Ideen beisteuern könntest. Dann könnte man ggf. auch mal über eine UDF nachdenken^^

  • Hi zusammen,
    habe nach einiger Zeit die Möglichkeiten von OpenCl dahingehend erweitert, nun auch mehrer Devices GLEICHZEITIG einen Kernel abarbeiten zu lassen.

    Die Umsetzung war garnicht so aufwendig, im Prinzip wird für jedes Device eine eigene CommandQueue erstellt, in welcher dann der Kernel abgearbeitet wird.
    Der benötigte Speicher wird einfach auf die Devices aufgeteilt. Bei GDI-Grafiken ist das recht unspektakulär, da der Speichertransfer meistens länger dauert wie das erstellen der Grafik. Die Grafik wird jetzt mit 2 Grafikkarten zwar doppelt so schnell erstellt, allerdings kostet der Overhead der Speicherverwaltung und Kopie in den Hauptspeicher so viel, dass die gesamte Aktion eher langsamer abläuft. Jedenfalls auf meinem Laptop mit AMD-APU, 2 Grafikchips und 1x 4-Kern-Prozessor, da hapert es etwas aufgrund des geshareten RAMs und der langsamen Transfers.
    Auf dem Desktop sieht das schon anders aus, da fällt der Hexacore-FX6300 nicht ganz so übel gegen die Grafikkarte Radeon 7790 ab.
    Auch die Speichertransfers "wuppen" da ganz anders.
    Aufgrund dieser teuern GDI-Speichertransfers machen aus o.g. Gründen aber leider mehrere Devices kaum einen Sinn bei GDI-Grafiken :wacko:

    Aber in Verbindung mit OpenGl zur Darstellung fallen diese Speichertransfers natürlich weg, ich vermute mal eine lineare Skalierung der Beschleunigung bei Verwendung von mehreren Devices. Das steht in Kürze an, incl. eines Testscriptes, bei dem man die vorhandenen Devices miteinander kombinieren kann.
    Leider habe ich weder einen Intel-Prozessor, noch eine Nvidia-Grafikkarte zum testen, und auch keine näheren Infos, wie gesharte Devices dort behandelt werden.

    Weiterhin bräuchte ich einen relativ aufwendigen Kernel, welcher nicht in Microsekunden abgearbeitet wird^^ Hat da jemand einen aufwendigen Grafikfilter oder sonst eine aufwendige Berechnung?
    Wie ich hier im Thread gesehen habe, sollten die bisherigen Scripte ja auch auf Intel/Nvidia laufen. Ist das so?

    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

    Einmal editiert, zuletzt von Andy (19. Oktober 2014 um 15:55)

  • auf Intel/Nvidia laufen. Ist das so?

    Kann ich nicht so bestätigen (980 oder 850M). Läuft auf Radeons, ältere und neue (5450 und R9 280X). Intel bisher nirgends. Komme aber gerade nicht an den Output ^^

  • Habe jetzt einige Varianten probiert, funktioniert hervorragend!
    Mir fehlt aber noch die letzte zündende Idee für die Funktionsgestaltung.
    Damit die Einfachheit der bisherigen Funktionen bei Berechnungen auf nur einem Device erhalten bleibt, habe ich Erweiterungen hinzugefügt.
    Bisher konnte man Platform und Device aussuchen über

    [autoit]

    _cl_GetDevice($device = "ALL", $num_device_used = 1, $num_platform_used = 1); explizit GPU oder CPU auswählen

    [/autoit]


    Alle diese Parameter waren bisher optional, bei _cl_GetDevice() wurde einfach das nächste ermittelte Device benutzt.
    Bei der expliziten Angabe von Platform und Device entsprechend das zu verwendende Gerät. Herauszufinden bspw. über das Testprogramm Devices64.au3
    Es wurde EINE CommandQueue erzeugt, auf die der Kernel zugriff.

    Intern habe ich nun die Funktion dahingehend erweitert, dass nun für JEDES Device eine eigene CommandQueue erzeugt wird, und nun auch explizit einem (zzt. noch gemeinsamen) Kernel zugewiesen werden kann. Wenn man das weitertreiben möchte, könnte man sogar verschiedene Kernel verwenden. Das würde dann dazu führen, gleichzeitig auf verschiedenen Devices unterschiedliche Kernel laufen zu lassen 8o Ich denke, da geht der Weg hin...

    Vorgehensweise:
    Ein bestimmtes Device soll benutzt werden, alle Funktionen nutzen wie gehabt!
    _CL_GetDevice("ALL", 2,1) ;zweites Device der ersten Platform wird benutzt
    Buffer erzeugen und Kernelparameter setzen wie gehabt
    _CL_RunKernel($DATA_SIZE , 0) ;gesamten Datenbereich berechnen, OpenCl legt die Workgroups selbst fest
    _CL_ReadBuffer($CL_buffer_out, $output_buffer) ;Puffer(alle Pixel) lesen


    Mehrere Devices sollen genutzt werden:
    Aufruf von _cl_GetDevice("MULTI") ;die weiteren Parameter werden ignoriert, für jede Platform und jedes Device wird eine eigene CommandQueue angelegt
    Buffer erzeugen wie gehabt
    Da die Kernel natürlich mit unterschiedlichen Argumenten gefüttert werden müssen (sonst berechnen alle Devices ja dasselbe Ergebnis) bekommt jeder Kernel seine eigenen Argumente, gleiche/konstante Argumente für alle Kernel wie bspw. Grafikgröße muss man natürlich nur einmal festlegen!

    Erstes Device:
    _CL_SetArg(0, "ptr*", $CL_buffer_out1) ;Grafikpufferadresse 1 für Grafikberechnung Kernel 1
    Dem Kernel wird ein Device (Command__Queue[Platform][Device]) zugewiesen
    _CL_RunKernel($DATA_SIZE, 0, $Comm_Queue[1][1]) ;Kernel berechnen auf Device Nr. 1 auf Platform Nr. 1

    Zweites Device:
    _CL_SetArg(0, "ptr*", $CL_buffer_out2) ;Grafikpufferadresse 2 für Grafikberechnung Kernel 2
    Dem Kernel wird ein Device (Command__Queue[Platform][Device]) zugewiesen
    _CL_RunKernel($DATA_SIZE, 0, $Comm_Queue[1][2]) ;Kernel berechnen auf Device Nr. 2 auf Platform Nr. 1

    Das wars schon^^
    Im Beispiel habe ich einen gemeinsamen Grafikpuffer verwendet, in den beide Devices an unterschiedlichen Adressen ihre Ergebnisse geschrieben haben.
    Man MUSS allerdings, da sich beim Auslesen/Schreiben des Puffers auf eine CommandQueue bezogen wird, diese auch angegeben werden.
    _CL_ReadBuffer($CL_buffer_out2, $output_buffer2, $comm_queue[1][2]);liest nur Puffer 2 aus, wenn Queue fertig berechnet
    _CL_ReadBuffer($CL_buffer_out, $output_buffer, $comm_queue[1][2]);Gesamten Puffer auslesen, wenn Queue fertig

    Der Witz liegt im "wenn Queue fertig". Ich habe die Funktionen der Einfachheit halber so gewrappert, dass der Puffer erst gelesen werden kann, wenn die Berechnung fertig ist!
    Bei unterschiedlich schnellen Devices macht es für die Weiterverarbeitung/Grafikdarstellung daher Sinn, auf das LANGSAMERE Device zu warten, oder direkt auf alle Berechnungen per _cl_Flush() zu warten.
    Man kann das im Beispiel mal testen, die Verwendung des schnelleren Device als Anhaltspunkt für das ReadBuffer führt zu Ruckeln im oberen Bildteil (wenn dort das langsamere Device "sitzt") :whistling: Da ja mit der Berechnung in der Schleife weitergemacht wird, erhält man so auch höhere FPS.
    Bei gleichzeitiger Berechnung an EINER Grafik bietet sich es also an, auf das jeweils langsamste Device zu warten. Hier suche ich ebenfalls noch eine Idee, bzw. empfehle den Flush.
    Logischerweise wird man, wie schon oben geschrieben, keine Frameverdoppelung bei Verwendung von 2 Grafikkarten feststellen. Das ReadBuffer() und anschließende Blitten dauert ca. 100x ( !!!!! ) länger als die Berechnung der Grafik durch den Kernel! Nutzwert also nur bei Verwendung von OpenGl oder SEHR aufwendigen Kerneln. (Irgendwo hab ich nen Mandelbulb-Kernel mit Schatten gesehen, der wäre geeignet^^)

    Ist die Vorgehensweise mit den Funktionserweiterungen so OK, oder hat jemand eine bessere Idee?

    Ich habe die "neuen" Funktionen und includes für diese Beispielanwendung mal umbenannt, damit ihr euch im Falle des Nichtfunktionierens nicht eure Installation zerschiesst.
    Das ist alles übel zusammengestrickt, wenn es aber so funktioniert, werde ich das "säubern"^^

    Als nächstes Vorhaben wird Devices64 neu erstellt, mit grafischer Ausgabe (Flug durchs Apfelmännchen) jedes einzelnen Device und einer Ausgabe mit Berechnung auf ALLEN Devices gleichzeitig!

    autoit.de/wcf/attachment/24843/

    Wer eine AMD-CPU und eine RADEON Grafikkarte hat, braucht nichts zu ändern, aber

    @Techmix und die anderen, welche MEHRERE Platforms (AMD/Intel/Nvidia)benutzen
    Zuerst in Zeile 67 der Tunnel_forum_all.au3 den Parameter $all=0 setzen und Script starten. Wenn das funktioniert (alte Funktionen werden verwendet, das erste erreichbare Device wird benutzt) bitte den Parameter $all=1 setzen.
    Schaut euch in der Konsole die verfügbaren Devices an (blaue Schrift)
    bspw.
    >Device verfügbar= 2;1;4;GPU;Bonaire;54581832;1551968152
    >Device verfügbar= 2;2;4;GPU;Bonaire;54581964;1551968152
    DAS ist die Platform
    DAS ist das Device
    Die CommandQueue müsst ihr nun in den Zeilen 104 bis 108 anpassen
    Also in Zeile 104 in $comm_queue[2][1]
    Zeile 107+108 $comm_queue[2][2]

    Ich habe keine Ahnung, ob das auch das "mischen" von verschiedenen Platformen funktioniert, man sollte zuerst bspw. die beiden Grafikkarten miteinander testen....
    Bei Fehlern bitte den Consolenanhang posten.
    Viel Erfolg 8o

    //EDIT
    noch eine kleine Anmerkung, wieso es in den allermeisten Fällen NICHTS bringt, die Kernel von einem Gespann CPU und GPU gleichzeitig berechnen zu lassen:
    Die Zeitersparnis errechnet sich aus dem Geschwindigkeitsfaktor GPU gegenüber CPU zu Zeitersparnis[%]= (1-Faktor/(1+Faktor))*100
    Mal angenommen die GPU ist 12x schneller als die CPU, dann erhält man für die Zeitersparnis gerade mal 7,7%. Das heisst, die GPU übernimmt 92,3% der Berechnungen, die CPU lediglich 7,7% der Berechnungen, unter der Vorraussetzung, CPU und GPU starten und beenden die Berechnungen gleichzeitig.
    Meistens ist jedoch der Faktor (die reine Kernelberechnungsszeit ohne Transfers) 20-50.
    Da kann sich jeder selbst ausrechnen, ob man, um 1-3% Geschwindigkeitsgewinn zu erzielen, den Aufwand mit der Berechnung auf einem Gespann GPU/CPU betreibt...

    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

    5 Mal editiert, zuletzt von Andy (21. Oktober 2014 um 08:31)

  • Moin!
    Habe jetzt deine neue Version getestet, hier erstmal die standart Ergebnisse:
    Cpu = 366fps, Gpu = 741fps, noSIMD(Cpu = 750fps, Gpu = 843fps)
    Wenn ich $all auf 1 setze:
    Multi = 1244fps, noSIMD(Multi = 1263) (ohne Änderungen)
    Multi = 1129fps, noSIMD(Multi = 1020) (mit Änderungen)
    Die Probleme sind wie folgt, der NoSimd-Kernel zeigt nur 25% des Bildschirminhaltes an und sobald ich Multi aktiviere habe ich ein schwarzes Bild...
    Das ich mit meinem i7 keine Verbesserung bei dem Simd-Kernel habe, hatte ich schon mal erwähnt, aber das fast die doppelten Fps OHNE Simd habe finde ich schon merkwürdig... Kann das daran liegen weil ich 8 Threads habe und der Kernel nur bis +3 zählt (convert_int4)??
    Auch finde ich die Idee der Multi-Kernel sehr gut, so kann man für jedes Device den gleichen Kernel oder eben unterschiedliche vergeben. Ich denke das wäre ein Vorteil, wenn man unterschiedliche Funktionsdurchläufe bei identischen Eingangsdaten durchführen muß - viele Fliegen mit einer Klappe :D
    Wenn ich die 'openCl_apfel_2d.au3' entsprechen der Multi-Version Anpasse, habe ich das gleiche schwarze Bild... Standart-Fps=384, Multi-Fps=588.

    SciTE Output


    Standart Output

    Multi ohne Änderungen

    Multi mit Änderungen

    Multi bei openCl_apfel_2d

    openCl_apfel_2d.au3
    [autoit]


    ;~ #include <opencl_easy.au3>
    #include <opencl_easy_all.au3>
    #include <GUIConstantsEx.au3>
    #include <WinAPI.au3>
    #include <GDIConstants.au3>
    #include <Misc.au3>

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

    Opt("GUIOnEventMode", 1)
    Opt("MouseCoordMode", 2)

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

    Global _
    $dll = DllOpen("user32.dll"), _
    $mouseflag = 1, _
    $fps = 0, _
    $NULL = 0, _
    $titeltext = "Apfelmännchen linke Maustaste reinzoomen / rechte Maustaste rauszoomen ", _
    $width, _ ; breite und höhe der gui/grafik, die Anzahl der "Threads" sollte ein Vielfaches von 16 sein!
    $height, _
    $hgui, _ ; GUI erstellen
    $hdc_gui, _ ; HDC holen zum blitten
    $DATA_SIZE, _ ; bildgröße=anzahl pixel (rgba)
    $hDC_bitmap, _
    $ptr_bitmap, _
    $hbmp_bitmap

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

    Global _
    $output_buffer, _ ; halleluja! bitmaps werden immer 16byte-aligned!
    $CL_buffer_out, _
    $maxiteration, _ ; maximale iterationen
    $centerx, _ ; mitte Bild in x-richtung
    $centery, _ ; mitte bild in y-richtung
    $apfelbreite, _ ; breite des Bildinhalts x-richtung
    $CL_DEBUGFLAG

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

    ;ich bin nicht der C-Freak, bitte also um Nachsicht mit meinen Kernels, Verbesserungsvorschläge werden selbstverständlich umgesetzt^^
    ;wer mag, kann den Kernel auch gerne als SIMD-Variante machen, das beschleunigt die Berechnung auf der CPU enorm!
    ;mit den farbfunktionen kann man endlos spielen^^
    Global _
    $KernelSource = _
    "__kernel void tunnelflug( __global int* output, const unsigned int width,const unsigned int height,const float centerx,const float centery,const uint maxiter,const float apfelbreite){" & @CRLF & _
    " uint iter;" & @CRLF & _ ;lokale variablen
    " float xtemp;" & @CRLF & _
    " uint threadid = get_global_id(0);" & @CRLF & _ ;thread-id, d.h. jedes pixel von 1 bis b*h
    " float py = threadid/height;" & @CRLF & _ ;pixelkoordinaten
    " float px = threadid%width;" & @CRLF & _ ;pixelkoordinaten
    " float x0 = (centerx-apfelbreite*0.5f)+px*apfelbreite/width;" & @CRLF & _ ;...Berechnungen
    " float y0 = (centery-apfelbreite*0.5f)+py*apfelbreite/height;" & @CRLF & _ ;...Berechnungen
    " float x = x0;" & @CRLF & _
    " float y = y0;" & @CRLF & _
    " for (iter=0;iter<maxiter;iter++){" & @CRLF & _ ;...Berechnungen
    " xtemp = x*x -y*y + x0;" & @CRLF & _
    " y = 2 * x * y + y0;" & @CRLF & _
    " x = xtemp;" & @CRLF & _
    " if ((x * x + y * y) > 4) break;" & @CRLF & _
    " }" & @CRLF & _
    " output[threadid] = 0x0008002* native_log2(native_log2((float) iter+1));" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t native_powr(iter,2)
    "}"

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

    ; ----------------------------------------------------------------------------------------------------------------------------------
    ; StartUP
    ; ----------------------------------------------------------------------------------------------------------------------------------

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

    Dim _
    $width = 256 * 3, _ ;breite und höhe der gui/grafik, die Anzahl der "Threads" sollte ein Vielfaches von 16 sein!
    $height = 256 * 3, _
    $hgui = GUICreate($titeltext, $width, $height, 1, 1), _ ;GUI erstellen
    $hdc_gui = _WinAPI_GetDC($hgui), _ ;HDC holen zum blitten
    $DATA_SIZE = $width * $height ;bildgröße=anzahl pixel (rgba)

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

    GUISetState()
    GUISetOnEvent($GUI_EVENT_CLOSE, "_Exit")
    GUISetOnEvent($GUI_EVENT_PRIMARYDOWN, "_mousedown")
    GUISetOnEvent($GUI_EVENT_SECONDARYDOWN, "_mousedown")
    AdlibRegister("_fps", 1000) ;FramesPerSecond

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

    ;bitmap erzeugen, in $ptr_bitmap steht nach dem Funktionsaufruf der Pointer auf die Pixeldaten
    $hDC_bitmap = _CreateNewBmp32($width, $height, $ptr_bitmap, $hbmp_bitmap) ;DC, Pointer auf die Bitmapdaten und ein Handle für GDI+....eine eierlegende Wollmilchsau

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

    Dim $all = 1
    ;ab hier gehts mit OpenCl los
    ;Zunächst werden die verfügbaren Geräte(Devices) gesucht.
    ;~ _CL_GetDevice("CPU") ;.. ALL oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    ;~ _CL_GetDevice("multi") ;.. ALL oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    If $all = 1 Then
    _CL_GetALLDevice("multi") ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    Else
    _CL_GetDevice("gpu", 1,1) ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    EndIf
    ;zweiter Parameter gibt bei z.b. mehreren Grafikkarten die Nummer des Geräts an

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

    ;Puffer für die Bitmap
    Dim _
    $output_buffer = DllStructCreate("dword[" & $DATA_SIZE & "]", $ptr_bitmap), _
    $CL_buffer_out = _CL_CreateBuffer($output_buffer), _
    $maxiteration = 89, _ ; maximale iterationen
    $centerx = -.5, _ ; mitte Bild in x-richtung
    $centery = 0, _ ; mitte bild in y-richtung
    $apfelbreite = 3, _ ; breite des Bildinhalts x-richtung
    $CL_DEBUGFLAG = 0 ; Debug-Ausgabe ausschalten, um Geschwindigkeit zu erhöhen

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

    $output_buffer1 = DllStructCreate("dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap);halleluja! bitmaps werden immer 16byte-aligned!
    $CL_buffer_out1 = _CL_CreateBuffer($output_buffer1)

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

    $output_buffer2 = DllStructCreate("dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap + $DATA_SIZE * 2);halleluja! bitmaps werden immer 16byte-aligned!
    $CL_buffer_out2 = _CL_CreateBuffer($output_buffer2)

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

    ;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*", $centerx)
    _CL_SetArg(4, "float*", $centery)
    _CL_SetArg(5, "uint*", $maxiteration)
    _CL_SetArg(6, "float*", $apfelbreite)

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

    ;Schleife zur Anzeige für den Flug
    While 1
    _steuerung()
    _CL_SetArg(3, "float*", $centerx)
    _CL_SetArg(4, "float*", $centery)
    _CL_SetArg(5, "uint*", $maxiteration)
    _CL_SetArg(6, "float*", $apfelbreite)
    ;~ _CL_RunKernelAll($DATA_SIZE / 2, 0) ; Kernel ausführen
    If $all = 1 Then
    _CL_SetArg(0, "ptr*", $CL_buffer_out1)
    _CL_RunKernel_ALL($DATA_SIZE / 2, 0, $comm_queue[2][1]) ;Kernel ausführeneue[1] & @crlf & '>Error code: ' & @error & @crlf) ;### Debug Console

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

    _CL_SetArg(0, "ptr*", $CL_buffer_out2)
    _CL_RunKernel_ALL($DATA_SIZE / 2, 0, $comm_queue[2][2]) ;Kernel ausführeneue[1] & @crlf & '>Error code: ' & @error & @crlf) ;### Debug Console
    _CL_ReadBuffer_ALL($CL_buffer_out2, $output_buffer2, $comm_queue[1][2]);Puffer(alle Pixel) lesen
    Else
    _CL_RunKernel($DATA_SIZE, 0) ; Kernel ausführen
    _CL_ReadBuffer($CL_buffer_out, $output_buffer) ; Puffer(alle Pixel) lesen
    EndIf
    _WinAPI_BitBlt($hdc_gui, 0, 0, $width, $height, $hDC_bitmap, 0, 0, $srccopy) ; Bitmap in die GUI blitten
    $fps += 1 ; Frames pro Sekunde zählen
    WEnd

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

    ;...das wars schon^^

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

    ; ----------------------------------------------------------------------------------------------------------------------------------
    ; Unterfunktionen
    ; ----------------------------------------------------------------------------------------------------------------------------------

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

    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, 1, DllStructGetSize($tBMI) - 4) ;Structgröße abzüglich der Daten für die Palette
    DllStructSetData($tBMI, 2, $iwidth)
    DllStructSetData($tBMI, 3, -$iheight) ;minus =standard = bottomup
    DllStructSetData($tBMI, 4, 1)
    DllStructSetData($tBMI, 5, 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ückgeben
    EndFunc ;==>_CreateNewBmp32

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

    Func _fps()
    WinSetTitle($hgui, "", $titeltext & $fps & " FPS")
    $fps = 0
    EndFunc

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

    Func _mousedown()
    $mouseflag = 1
    EndFunc

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

    Func _steuerung()
    If WinActive("Apfelmännchen") Then
    If _IsPressed("01", $dll) Then _leftmouse()
    If _IsPressed("02", $dll) Then _rightmouse()
    EndIf
    EndFunc

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

    Func _leftmouse() ;zoom in
    If $mouseflag Then ;neues center
    $posx = MouseGetPos(0)
    $posy = MouseGetPos(1)
    $centerx = ($centerx - $apfelbreite / 2) + $posx * $apfelbreite / $width
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $centerx = ' & $centerx & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $centery = ($centery - $apfelbreite / 2) + $posy * $apfelbreite / $height
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $centery = ' & $centery & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $mouseflag = 0
    EndIf
    $maxiteration /= 0.99785 ;Iterationen erhöhen
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $maxiteration = ' & $maxiteration & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $apfelbreite *= 0.985
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $apfelbreite = ' & $apfelbreite & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    EndFunc

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

    Func _rightmouse() ;zoom out
    If $mouseflag Then ;neues center
    $posx = MouseGetPos(0)
    $posy = MouseGetPos(1)
    $centerx = ($centerx - $apfelbreite / 2) + $posx * $apfelbreite / $width
    $centery = ($centery - $apfelbreite / 2) + $posy * $apfelbreite / $height
    $mouseflag = 0
    EndIf
    $maxiteration *= 0.99785 ;Iterationen erhöhen
    $apfelbreite /= 0.985
    EndFunc

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

    Func _Exit()
    Exit
    EndFunc ;==>_Exit

    [/autoit]

    Grüße!

  • _CL_SetArg(0, "ptr*", $CL_buffer_out2)
    _CL_RunKernel_ALL($DATA_SIZE / 2, 0, $comm_queue[2][2]) ;Kernel ausführeneue[1] & @crlf & '>Error code: ' & @error & @crlf) ;### Debug Console
    _CL_ReadBuffer_ALL($CL_buffer_out2, $output_buffer2, $comm_queue[1][2]);Puffer(alle Pixel) lesen

    Schau mal, bei ReadBuffer die $comm_queue sollte [2][1] sein!

    Die Probleme sind wie folgt, der NoSimd-Kernel zeigt nur 25% des Bildschirminhaltes an und sobald ich Multi aktiviere habe ich ein schwarzes Bild...


    Die 25% werden deshalb dargestellt, weil du auch nur 1/4 der Daten berechnest. Schau mal in der Zeile mit RunKernel() da müsste die Datasize ohne /4 rein.

    Bei Multi schwarzes Bild, das heisst, der Grafikbuffer wird entweder nicht richtig ausgelesen (setz mal ALLE $comm_queue auf denselben Wert) oder beschrieben.
    Du kannst auch, um überhaupt mal zu testen ob Multi funktioniert, nur EIN Device rechnen lassen (alle $comm_queue auf denselben Wert und nur einen buffer)

    Dein Apfelmännchen-script läuft bei mir. Ich sehe ein halbes Apfelmännchen, was sich auch zoomen lässt. Ich habe lediglich die 3 Zeilen mit den $comm_queue an meine Werte angepasst. Man sieht natürlich nur die obere/untere Hälfte, du hast einen Logikfehler drin.
    Du erstellst richtigerweise zwei einzelne Grafikpuffer, allerdings müssen diese natürlich zusammengefasst werden, ich habe im Tunnel-Script dazu einen dritten Grafikpuffer an dieser Position erstellt, der so groß ist wie beide zusammen.
    Damit die getrennte Berechnung funktioniert, musst du auch beide Teile der Grafik getrennt berechnen! Das beinhaltet so gut wie alle Kernel-Parameter.
    Oberer Bildteil für sich, und unterer Bildteil für sich. Als ob du die Grafik in ein eigenes Fenster zeichnen würdest!

    Das ich mit meinem i7 keine Verbesserung bei dem Simd-Kernel habe, hatte ich schon mal erwähnt, aber das fast die doppelten Fps OHNE Simd habe finde ich schon merkwürdig..

    Ich nicht^^, bei schwarzem Bildschirm findet kein Transfer von "Daten" statt. So clever ist GDI seit den 90ern^^
    Du bekommst also nicht die FPS angezeigt, sondern die Schleifendurchläufe pro Sekunde, bitblt macht also nix.

  • Mahlzeit,
    Dieses Script war noch von morgens als ich meinen Post nicht absenden konnte - war nicht aktuell... Ich hatte bereits alle Möglichkeiten bei $comm_queue durchprobiert, keine Änderungen immer ein schwarzes Bild!

    Und die Simd Test hatte ich ohne Multi, mit Cpu oder Gpu, durchgeführt. Ich verstehe einfach nicht wieso beide Devices ein ähnliches Verhalten bei den Simd-Kernel auweisen, also langsammer als bei dem normalen Kernel sind. Kann das an der OpenCL Installation liegen? Zu alt wird meine Hardware auch nicht sein, und wenn ich OpenCL z.B. in Mencoder verwende habe ich auch einen Performance plus - OK, ich weiß ja nicht wie deren Kernel aufgebaut ist... Aber grundsätzlich gesehen. Auch scheint Cuda sehr gut zu laufen, habe damit bei DirectShow Encoding 250% Leistungszuwachs! Aber dieser Simd-Kernel...

    Grüße!

  • Ich verstehe einfach nicht wieso beide Devices ein ähnliches Verhalten bei den Simd-Kernel auweisen, also langsammer als bei dem normalen Kernel sind

    Wenn du die Laufzeiten der einzelnen Kernel vergleichen willst, musst du dir die Kernellaufzeiten zurückgeben lassen!

    [autoit]

    _CL_RunKernel($DATA_SIZE, 0) ;Kernel ausführen,
    $kerneltime = profile_cl_event()

    [/autoit]

    Also funktionierendes Script erstellen (incl. angezeigter Grafik! ) und DANN die realen Laufzeiten der Kernel vergleichen. Dazu würde ich die Schleife ca. 1000x durchlaufen lassen.

    Diese FPS-Anzeige ist, wie oben erläutert, nur funktionabel wenn auch eine Grafik angezeigt wird, und bei "schwarzem Bildschirm" völlig nutzlos!

    Ich hatte bereits alle Möglichkeiten bei $comm_queue durchprobiert, keine Änderungen immer ein schwarzes Bild!


    Also der "Normalmodus" mit $all=0 funktioniert, nur bei $all=1 nicht?
    OK, erstelle mal EINEN Buffer und lass das Script mit einem Device, einem Kernel und einem ReadBuffer laufen, mit $command_queue[1][1] (sollte bei dir die CPU sein). Funktioniert das mit Anzeige?
    Wenn ja, teste das mit allen Devices EINZELN ( also nur in [2][1] (GPU1) und [2][2] (GPU2) ändern!).
    Wenn das nicht funktioniert, poste mal dein Script! Ich vermute immer noch, dass du einen Logikfehler bei den Buffern hast.

    Es gibt keine OpenCL (und auch keine CUDA)-Funktion "verteil mal diesen Kernel auf alle Devices, mach ne Lastanalyse, pass die Kernelargumente an und alles ist schick!"
    Woher soll auch die API wissen, an welchem der Kernelargumente man etwas verändern muss, damit nur ein bestimmter Teil bespw. einer Grafik berechnet werden soll? Wenn du 4 Devices hast, bei denen nach der Lastanalyse die jeweilige Belastung 10%/25%/25%/40% ist, woher soll dei API wissen, dass du 4 einzelne Buffer der jeweiligen Größe brauchst, und natürlich auch 4 Start-Pointer innerhalb der Grafik, an denen die Berechnungen anfangen?

    Das bringt mich aber auf eine Idee!
    Du solltest ggf. direkt 4 Grafiken erstellen (schon mit _CreateNewBmp32()) diese in 4 Buffer überführen, jeweils einzeln berechnen/auslesen und nachher diese 4 Teile in EINE GUI blitten.
    Ich denke, dieser Ansatz beseitigt dein Logikproblem.
    Das Problem bei diesem Ansatz ist, dass durch 4x ReadBuffer die FPS nur noch 1/4 sind....^^

    Bei deinem Beispielcode hast du nur 2 Buffer, den oberen und den unteren Teil. Frag dich mal, wie BitBlt den oberen zusammen mit dem unteren Teil der Grafik in die GUI blitten soll, wenn du per ReadBuffer nur eine Hälfte ausliest?

    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

    Einmal editiert, zuletzt von Andy (22. Oktober 2014 um 14:30)

  • Hy Andy,
    habe jetzt ein Bild bekommen! Lag daran, das ich vergaß _CL_ReadBuffer_ALL() anzupassen - glatt überlesen...
    ABER es gibt noch immer ein Problem, und zwar wird mir immer nur ein Bildausschnitt angezeigt. Ist schon komisch, habe aber selber noch getüftelt und dabei ist folgendes rausgekommen:

    [autoit]


    $all = 1

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

    ; ...

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

    $output_buffer1 = DllStructCreate("dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap) ; Puffer 1
    $CL_buffer_out1 = _CL_CreateBuffer($output_buffer1)

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

    $output_buffer2 = DllStructCreate("dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap + $DATA_SIZE * 2) ; Puffer 2
    $CL_buffer_out2 = _CL_CreateBuffer($output_buffer2)

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

    ; ...

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

    While 1 ;endless loop...
    $steps += $step ;schrittweite für den nächsten frame erhöhen...
    _CL_SetArg(3, "float*", $steps) ;...und an kernel übergeben

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

    _CL_SetArg(0, "ptr*", $CL_buffer_out1)
    _CL_RunKernel_ALL($DATA_SIZE / 8, 0, $comm_queue[2][1]) ; Kernel ausführen GPU 1
    _CL_SetArg(0, "ptr*", $CL_buffer_out2)
    _CL_RunKernel_ALL($DATA_SIZE / 8, 0, $comm_queue[2][2]) ; Kernel ausführen GPU 2

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

    _CL_ReadBuffer_ALL($CL_buffer_out1, $output_buffer1, $comm_queue[2][1]) ; Puffer 1 lesen
    _CL_ReadBuffer_ALL($CL_buffer_out2, $output_buffer2, $comm_queue[2][2]) ; Puffer 2 lesen

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

    _WinAPI_BitBlt($hdc_gui, 0, 0, $width, $height, $hDC_bitmap, 0, 0, $srccopy) ;Bitmap in die GUI blitten
    $fps += 1 ;Frames pro Sekunde zählen
    WEnd

    [/autoit]

    Jetzt bekomme ich 2 Halbbilder, jedes in seiner richtigen Position aber beide zeichnen nur den oberen Bildinhalt. Lese ich den Puffer 1 nicht aus (wie in deinem Beispiel) wird nur der untere Bildteil gezeichnet.
    Testläufe: "GPU" = 780Fps (richtiges Bild), "MULTI" = 580Fps (Doppelbildausschnitt)
    Und habe bei mir festgestellt das die 64Bit Ausführung bei allen Tests mit allen Kerneln 40 bis 50 Fps mehr bringt!

    Was anderes...
    Du weißt ja ich habe großes Interesse an OpenCL, Und wenn ich mir deine 'UDF' mal anschaue denke ich das vielleicht einige Änderungen das Arbeiten erleichtern könnten!

    _CL_SetKernel($sKernel, $aCommand_Queue)
    Fehlt noch völlig... Man könnte den zweiten Parameter so abfragen, wenn ein Array übergeben ist wird der Kernel auf alle Devices verteilt und wenn ein $command_queue angegeben ist, soll der Kernel eben nur auf den einem Device ausgeführt werden. ...Diese globale Variable $KernelSource ist völlig überflüssig und verwirrt den Techmix nur ;)

    [Überholt...]
    _CL_GetALLDevice("MULTI")
    Ich weiß nicht ob es hier noch einen Sinn macht die weiteren Parameter bei zu behalten. Was hier noch ein sinniger Parameter wäre ist die '$num_device_used',
    wenn man die auf -1 stellt könntest du noch einen "MIX" Betrieb anbieten welcher auch noch die CPU einschließt. Ja, ich habe dein Post gelesen und bin auch deiner
    Meinung das ein Mix-Betrieb, begründet duch die CPU-Leitung, wenig Sinn ergibt. Aber sollte jemand einen Video-Encoding-Ultra-Kernel verarbeiten wollen ist er
    sicherlich Dankbar für das quentchen Mehrleistung. Auch würden die Funktionen einfacher zu verstehen sein wenn du das z.B so $command_queue = _CL_GetALLDevice()
    lösen würdest.

    _CL_CreateBuffer($output_buffer)
    Immer diese manuelle Zuweisung...

    _CL_RunKernel_ALL($DATA_SIZE / 8, 0, $command_queue[2][1])
    Bau die Funktion so auf, das diese intern die $command_queue´s verwaltet wenn diese als Array übergeben wurden.
    P.S: Wenn ich das mache _CL_RunKernel_ALL($DATA_SIZE / 4, ...) gibts bei mir nen übelen Grafiktreiberabsturz ;)

    _CL_ReadBuffer_ALL($CL_buffer, $read_buffer, $Queue = $command_queue)
    Wieso muß ich noch den $command_queue angeben obwohl ich doch bereits entsprechende Buffer angegeben habe?? Lasse ich diesen aus, schwarzes Bild.
    [/Überholt...]

    ...

    Oki direkt angegangen und umgesetzt :D
    _cl_GetALLDevice2($device, ByRef $aOutput_Buffer, ByRef $aCommand_Queue, $iDeviceNr=0)
    Es können alle Arten von Devices verarbeitet werden (cpu, gpu, multi).
    $aCommand_Queue und $aOutput_Buffer werden gleich in entsprechender Größe für die jeweilige Platform erstellt.
    Die $iDeviceNr bestimmt im Single Modus mit welchem Device gearbeitet wird.
    Ausserdem werden jetzt die Platformen, Queues und Buffer automatisch zugeordnet, der Anwender wird jetzt weniger verwirrt ;)

    _CL_RunKernel_ALL2($Global, $Local, $aCommand_Queue, $aOutput_Buffer="")
    Es wird das Array $aCommand_Queue durchgearbeitet, immer wird erst _CL_SetArg() und anschliessend _CL_RunKernel_ALL() ausgeführt.
    $Global und $Local werden wie zuvor behandelt.
    Wenn $aOutput_Buffer leer ist, wird die ursprüngliche Funktion durchlaufen. <- Also ohne Loop und Array´s.

    _CL_CreateBuffer_ALL2(ByRef $aOutput_Buffer, $sStructCreate1,$sPointer1, $sStructCreate2="",$sPointer2="", ...)
    Hier habe ich keine andere Lösung gefunden, als alle möglichen Argumente optional anzuhängen ($sStructCreateXX,$sPointerXX).
    Das Array $aOutput_Buffer wird solange durchlaufen, wie Argumente vorhanden sind. So werden alle Buffer mit einem Befehl erstellt.

    _CL_ReadBuffer_ALL2($aOutput_Buffer, $aCommand_Queue, $command_queue="")
    Es wird das Array $aCommand_Queue durchgearbeitet, es wird immer ein _CL_ReadBuffer_ALL() ausgeführt.
    Wenn $command_queue nicht leer ist, wird die ursprüngliche Funktion durchlaufen. <- Also ohne Loop und Array´s.

    Aufrufe
    [autoit]


    ; ...
    ; ...
    Local $aOutput_Buffer, $aCommand_Queue, $iDeviceNr, $iDataSize
    Dim $all = 1
    If $all = 1 Then
    $iDataSize = $DATA_SIZE / 8
    _CL_GetALLDevice2("MULTI", $aOutput_Buffer, $aCommand_Queue)
    _CL_CreateBuffer_ALL2($aOutput_Buffer, "dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap, "dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap + $DATA_SIZE * 2)
    Else
    $iDeviceNr = 0
    $iDataSize = $DATA_SIZE / 4
    _CL_GetALLDevice2("gpu", $aOutput_Buffer, $aCommand_Queue, $iDeviceNr)
    _CL_CreateBuffer_ALL2($aOutput_Buffer, "dword[" & $DATA_SIZE & "]", $ptr_bitmap)
    EndIf
    ; ...
    ; ...
    While 1 ;endless loop...
    $steps += $step ;schrittweite für den nächsten frame erhöhen...
    _CL_SetArg(3, "float*", $steps) ;...und an kernel übergeben

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

    _CL_RunKernel_ALL2($iDataSize, 0, $aCommand_Queue, $aOutput_Buffer) ; Kernel ausführen, egal ob Multi oder Single...
    _CL_ReadBuffer_ALL2($aOutput_Buffer, $aCommand_Queue) ; Puffer auslesen

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

    _WinAPI_BitBlt($hdc_gui, 0, 0, $width, $height, $hDC_bitmap, 0, 0, $srccopy) ;Bitmap in die GUI blitten
    $fps += 1 ;Frames pro Sekunde zählen
    WEnd

    [/autoit]

    Die neuen Funktionen laufen mit Multi und Single ohne das man große Änderungen durchführen muß, lediglich die $iDeviceNr muß angepasst (oder ausgelassen) werden.
    In der Funktion _cl_GetALLDevice2() gibt es allerdings noch Schwierigkeiten, und zwar bekomme ich die CPU micht zum laufen...
    Ich habe ich an der $platform_device die $comm_queue[$i][$t] angehängt - um die PID´s mit denen in $comm_queue abzugleichen... Und es Funktioniert!
    Aber das beste ist die automatische Zuordnung von Platformen, Queue´s und Buffern! Auch das der Loop jetzt übersichtlicher und vor allem flexibler ist empfinde ich als angenehm.
    Aber du mußt noch unbedingt die Funktion _CL_SetKernel() erstellen bzw. entkoppeln aus _cl_GetALLDevice()...

    Und langsamm verstehe ich was du schriebst, im Multi-Modus brauchst du auch einen angepassten Kernel! Mit nur einem Kernel (wie jetzt) bleibt nur die Möglichkeit diesen 'Mono' Kernel mit _CL_SetArg() entsprechende Start und Stopp Positionen zu übergeben - wenn natürlich der Kernel entsprechend umgeschrieben wurde :D
    Ich habe bei mir jetzt schon bestimmt 36 GraKa Abstürze provoziert, habe noch Logikprobleme *Grins* betreffend der Buffergröße, Start und Stop Zuweisung einer Multi-Kernel Modifikation des Tunnel-Kernel (NoSimd)... Kernelintern wird ja width*height berechnet, entspricht also $DATA_SIZE, die Rückgabe des Kernel soll ja nur $DATA_SIZE/2 groß sein - entsprechend müssen die Buffer also auch nur die Größe von $DATA_SIZE/2 haben. Richtig? ;)
    Ich weiß jetzt aber wieso ich vorher andauernde Abstürze in dem anderen Script hatte: Ich hatte 2 Buffer voller Größe erstellt, und bei der Rückgabe des 2ten Buffers wurde fleißig in den Speicherbereich 'von anderswo' reingeschrieben, weil der Buffer ja um 50% zu groß ist ;) <- HA, mache Fortschritte! Ich Programmiere ja nur Offline - und da ist das mit den schnellen Lösungen Googlen etwas schwierig, da kommts auf gute Vorbereitung und reichlich Griebenschmalz an :D !Und somit ist das Doppelbild-Problem gar kein Problem gewesen, sondern genau so sollte es auch sein!
    ...Habe jetzt einige Stunden mit den eigenen Funktionen verbracht, und muß sagen das diese nur einen Sinn ergeben wenn die Kernel-Parameter wenigstens an der ersten Stelle (output) immer identisch sind.
    Aber ich habe Multi Versionen des Kernel erstellen können *Yipie* welche aber noch optimierungsbedürftig sind!

    MultiKernel
    [autoit]


    ; Ich komme einfach nicht darauf WIE ich den Kernel sagen kann, das er bei 'stop' aufhören soll.. (if (threadid = stop+1) {break}) ?!?
    ; Momentan wird dieser trotzdem durchlaufen aber alle Berechnungen und die Rückgabe werden ausgelassen
    $KernelSource = _ ; Multi NoSimd-Kernel
    "__kernel void tunnelflug(__global int* output, const unsigned int width, const unsigned int height, const float stepx, const int start, const int stop)" & @CRLF & _
    "{" & @CRLF & _
    " uint threadid = get_global_id(0) + start - 1;" & @CRLF & _ ;thread-id, d.h. jedes pixel von 1 + start bis b*h
    " if (threadid < stop + 1) {" & @CRLF & _
    " 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 - start + 1] = 0xFF000000+(t<<16)+(t<<8)+t;" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t (- start)
    " }" & @CRLF & _
    "}"

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

    ; ...

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

    ; So setze ich die Argumente entsprechend der Gpu: Gpu1=Obere Bildhälfte, Gpu2=Untere Bildhälfte
    ; ...Ich muß dazu gestehen das ich mich mit dem Verstehen von Grafikalgorythmen schwer tue, mir entzieht sich die Vorgehensweise und vor allem bleibt der 'Aha' Effekt aus.
    if $iGPU = 1 then
    _CL_SetArg(4, "uint*", $iGPU) ; Start ;; $iGPU soll einfach eine 1 darstellen...
    _CL_SetArg(5, "uint*", $DATA_SIZE / 2) ; Ende
    Else
    _CL_SetArg(4, "uint*", $DATA_SIZE / 2) ; Start
    _CL_SetArg(5, "uint*", $DATA_SIZE) ; Ende
    EndIf
    _CL_RunKernel_ALL()

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

    ; ...

    [/autoit]

    Ich hänge mal die 'tunnel_forum_ALL2.au3' an, da kannst du die Funktionen und Kernel testen.

    Grüße!

    P.S:
    Wir sollten uns mal einige alternative Sczenarios für den Einsatz für OpenCL überlegen, um vielleicht einen produktiven Einsatz von OCL in Au3 aufzuzeigen. Vielleicht könnte man ja mit Replace-Funktionen anfangen (bringen in ASM auch mächtig Speed) und es später vielleicht bis zu den _Crypt Funktionen ausbauen.
    Diese Grafik-Kernel sind ja ganz nett, aber einen richtigen Nutzen zeigen diese den Anwender ja nicht auf; Außer eben das Speed+ auf einigen Algoryhtmen...

    P.P.S:
    Ich habe mir eine Master-Arbeit aus dem Netz geladen welche glücklicherweise auf Deutsch verfasst ist und sich um das Thema OpenCL dreht. Da habe jetzt folgende interessante Info rauslesen können und zwar gibt es in OCL 2 Modi´s (?), einmal die Paralesierung und einmal das Threding! O, den ersten Punkt verwenden wir gerade weil wir Mathematisch vorgehen, das Threading aber bietet sich doch z.B. für Stringbearbeitung an!
    Kann ich mir das so vorstellen, wie man es aus C oder FreeBasic kennt?? Ich Denke da benötigt man ein OpenCL-Programm mit entsprechenden Kerneln welche dann die Threads darstellen...

  • Hi,

    Testläufe: "GPU" = 780Fps (richtiges Bild), "MULTI" = 580Fps (Doppelbildausschnitt)

    yepp, hängt daran, dass du 2x den Puffer von der Graka ins RAM transferierst, was, wie schon oben beschrieben, wesentlich länger dauert als die eigentliche Berechnung!
    Die oberen Halbbilder siehst du deshalb, weil du diese jeweils auch genau so darstellst. Lt. deinem Script ist das auch so gewünscht :P

    Fehlt noch völlig... Man könnte den zweiten Parameter so abfragen, wenn ein Array übergeben ist wird der Kernel auf alle Devices verteilt und wenn ein $command_queue angegeben ist, soll der Kernel eben nur auf den einem Device ausgeführt werden. ...Diese globale Variable $KernelSource ist völlig überflüssig und verwirrt den Techmix nu


    Den Kernel auf alle Devices zu verteilen macht imho ja nur dann Sinn, wenn diese Devices auch verwendet werden sollen. Wir könnten das nun natürlich als "Standard" implementieren, also einfach erstmal per se den Kernel auf ALLE verfügbaren Devices verteilen (wird ja aber schon gemacht, über die einzelnen Command-Queues)

    So wie es aussieht, funktioniert ja bei dir schonmal die Darstellung so, wie ich mir das vorgestellt hatte. :thumbup:

    Zitat

    ; Ich komme einfach nicht darauf WIE ich den Kernel sagen kann, das er bei 'stop' aufhören soll.. (if (threadid = stop+1) {break}) ?!?
    ; Momentan wird dieser trotzdem durchlaufen aber alle Berechnungen und die Rückgabe werden ausgelassen

    Naja, ich denke, es gibt noch Erklärungsbedarf....
    Du musst, ausser bei speziellen Aufgaben (wenn Ergebnisse synchronisiert werden müssen), den Kernel überhaupt nicht bearbeiten!
    Durch die Kernelparameter legst du fest, welche "Teilbereiche" des Bildes berechnet werden.
    Beispiel Apfelmännchen/Mandelbrotmenge.
    Welchen "Bildausschnitt" du angezeigt bekommst, bestimmen AUSSCHLIESSLICH die Kernelparameter. Der Kernel (Rechenvorschrift) berechnet immer nur EIN EINZIGES Pixel (ThreadID)!
    OpenCL verteilt nun diese eine Rechenvorschrift auf alle verfügbaren Recheneinheiten.
    Ich schreib mal ein Beispiel

    Spoiler anzeigen
    [autoit]

    ;~ #include <opencl_easy.au3>
    #include <opencl_easy_all.au3>
    #include <GUIConstantsEx.au3>
    #include <WinAPI.au3>
    #include <GDIConstants.au3>
    #include <Misc.au3>

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

    Opt("GUIOnEventMode", 1)
    Opt("MouseCoordMode", 2)

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

    Global _
    $dll = DllOpen("user32.dll"), _
    $mouseflag = 1, _
    $fps = 0, _
    $NULL = 0, _
    $titeltext = "Apfelmännchen linke Maustaste reinzoomen / rechte Maustaste rauszoomen ", _
    $width, _ ; breite und höhe der gui/grafik, die Anzahl der "Threads" sollte ein Vielfaches von 16 sein!
    $height, _
    $hgui, _ ; GUI erstellen
    $hdc_gui, _ ; HDC holen zum blitten
    $DATA_SIZE, _ ; bildgröße=anzahl pixel (rgba)
    $hDC_bitmap, _
    $ptr_bitmap, _
    $hbmp_bitmap

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

    Global _
    $output_buffer, _ ; halleluja! bitmaps werden immer 16byte-aligned!
    $CL_buffer_out, _
    $maxiteration, _ ; maximale iterationen
    $centerx, _ ; mitte Bild in x-richtung
    $centery, _ ; mitte bild in y-richtung
    $apfelbreite, _ ; breite des Bildinhalts x-richtung
    $CL_DEBUGFLAG

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

    ;ich bin nicht der C-Freak, bitte also um Nachsicht mit meinen Kernels, Verbesserungsvorschläge werden selbstverständlich umgesetzt^^
    ;wer mag, kann den Kernel auch gerne als SIMD-Variante machen, das beschleunigt die Berechnung auf der CPU enorm!
    ;mit den farbfunktionen kann man endlos spielen^^
    Global _
    $KernelSource = _
    "__kernel void tunnelflug( __global int* output, const unsigned int width,const unsigned int height,const float centerx,const float centery,const uint maxiter,const float apfelbreite){" & @CRLF & _
    " uint iter;" & @CRLF & _ ;lokale variablen
    " float xtemp;" & @CRLF & _
    " uint threadid = get_global_id(0);" & @CRLF & _ ;thread-id, d.h. jedes pixel von 1 bis b*h
    " float py = threadid/height;" & @CRLF & _ ;pixelkoordinaten
    " float px = threadid%width;" & @CRLF & _ ;pixelkoordinaten
    " float x0 = (centerx-apfelbreite*0.5f)+px*apfelbreite/width;" & @CRLF & _ ;...Berechnungen
    " float y0 = (centery-apfelbreite*0.5f)+py*apfelbreite/height;" & @CRLF & _ ;...Berechnungen
    " float x = x0;" & @CRLF & _
    " float y = y0;" & @CRLF & _
    " for (iter=0;iter<maxiter;iter++){" & @CRLF & _ ;...Berechnungen
    " xtemp = x*x -y*y + x0;" & @CRLF & _
    " y = 2 * x * y + y0;" & @CRLF & _
    " x = xtemp;" & @CRLF & _
    " if ((x * x + y * y) > 4) break;" & @CRLF & _
    " }" & @CRLF & _
    " output[threadid] = 0x0008002* native_log2(native_log2((float) iter+1));" & @CRLF & _ ;Pixel schreiben 0xFFBBGGRR, BB=GG=RR=t native_powr(iter,2)
    "}"

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

    ; ----------------------------------------------------------------------------------------------------------------------------------
    ; StartUP
    ; ----------------------------------------------------------------------------------------------------------------------------------

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

    Dim _
    $width = 256 * 3, _ ;breite und höhe der gui/grafik, die Anzahl der "Threads" sollte ein Vielfaches von 16 sein!
    $height = 256 * 3, _
    $hgui = GUICreate($titeltext, $width, $height, 1, 1), _ ;GUI erstellen
    $hdc_gui = _WinAPI_GetDC($hgui), _ ;HDC holen zum blitten
    $DATA_SIZE = $width * $height ;bildgröße=anzahl pixel (rgba)

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

    GUISetState()
    GUISetOnEvent($GUI_EVENT_CLOSE, "_Exit")
    GUISetOnEvent($GUI_EVENT_PRIMARYDOWN, "_mousedown")
    GUISetOnEvent($GUI_EVENT_SECONDARYDOWN, "_mousedown")
    AdlibRegister("_fps", 1000) ;FramesPerSecond

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

    ;bitmap erzeugen, in $ptr_bitmap steht nach dem Funktionsaufruf der Pointer auf die Pixeldaten
    $hDC_bitmap = _CreateNewBmp32($width, $height, $ptr_bitmap, $hbmp_bitmap) ;DC, Pointer auf die Bitmapdaten und ein Handle für GDI+....eine eierlegende Wollmilchsau

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

    Dim $all = 1
    ;ab hier gehts mit OpenCl los
    ;Zunächst werden die verfügbaren Geräte(Devices) gesucht.
    ;~ _CL_GetDevice("CPU") ;.. ALL oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    ;~ _CL_GetDevice("multi") ;.. ALL oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    If $all = 1 Then
    _CL_GetALLDevice("multi") ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    Else
    _CL_GetDevice("gpu", 1,1) ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.
    EndIf
    ;zweiter Parameter gibt bei z.b. mehreren Grafikkarten die Nummer des Geräts an

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

    ;Puffer für die Bitmap
    Dim _
    $output_buffer = DllStructCreate("dword[" & $DATA_SIZE & "]", $ptr_bitmap), _
    $CL_buffer_out = _CL_CreateBuffer($output_buffer), _
    $maxiteration = 89, _ ; maximale iterationen
    $centerx = -.5, _ ; mitte Bild in x-richtung
    $centery = 0, _ ; mitte bild in y-richtung
    $apfelbreite = 3, _ ; breite des Bildinhalts x-richtung
    $CL_DEBUGFLAG = 0 ; Debug-Ausgabe ausschalten, um Geschwindigkeit zu erhöhen

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

    $output_buffer1 = DllStructCreate("dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap);halleluja! bitmaps werden immer 16byte-aligned!
    $CL_buffer_out1 = _CL_CreateBuffer($output_buffer1)

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

    $output_buffer2 = DllStructCreate("dword[" & $DATA_SIZE / 2 & "]", $ptr_bitmap + $DATA_SIZE * 2);halleluja! bitmaps werden immer 16byte-aligned!
    $CL_buffer_out2 = _CL_CreateBuffer($output_buffer2)

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

    $output_buffer = DllStructCreate("dword[" & $DATA_SIZE & "]", $ptr_bitmap);halleluja! bitmaps werden immer 16byte-aligned!
    $CL_buffer_out = _CL_CreateBuffer($output_buffer)

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

    ;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*", $centerx)
    _CL_SetArg(4, "float*", $centery)
    _CL_SetArg(5, "uint*", $maxiteration)
    _CL_SetArg(6, "float*", $apfelbreite)

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

    ;Schleife zur Anzeige für den Flug
    While 1
    _steuerung()
    _CL_SetArg(3, "float*", $centerx)
    _CL_SetArg(5, "uint*", $maxiteration)
    _CL_SetArg(6, "float*", $apfelbreite)
    ;~ _CL_RunKernelAll($DATA_SIZE / 2, 0) ; Kernel ausführen
    If $all = 1 Then
    ;obere Hälfte Bild
    _CL_SetArg(4, "float*", $centery);Die Mitte des Bildes -> oberer Teil
    _CL_SetArg(0, "ptr*", $CL_buffer_out1)
    _CL_RunKernel_ALL($DATA_SIZE / 2, 0, $comm_queue[1][1]) ;Datasize /2 nur die Hälfte ALLER Pixel also nur OBERER Teil
    _CL_ReadBuffer_ALL($CL_buffer_out, $output_buffer, $comm_queue[1][1]);Puffer(alle oberen Pixel) lesen

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

    ;untere Hälfte Bild
    _CL_SetArg(4, "float*", $centery+$apfelbreite/2);unterer Teil des Bildes
    _CL_SetArg(0, "ptr*", $CL_buffer_out2)
    _CL_RunKernel_ALL($DATA_SIZE / 2, 0, $comm_queue[1][2]) ;Kernel ausführeneue[1] & @crlf & '>Error code: ' & @error & @crlf) ;### Debug Console
    _CL_ReadBuffer_ALL($CL_buffer_out, $output_buffer, $comm_queue[1][2]);Puffer(alle Pixel) lesen
    Else
    _CL_RunKernel($DATA_SIZE, 0) ; Kernel ausführen
    _CL_ReadBuffer($CL_buffer_out, $output_buffer) ; Puffer(alle Pixel) lesen
    EndIf
    _WinAPI_BitBlt($hdc_gui, 0, 0, $width, $height, $hDC_bitmap, 0, 0, $srccopy) ; Bitmap in die GUI blitten
    $fps += 1 ; Frames pro Sekunde zählen
    WEnd

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

    ;...das wars schon^^

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

    ; ----------------------------------------------------------------------------------------------------------------------------------
    ; Unterfunktionen
    ; ----------------------------------------------------------------------------------------------------------------------------------

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

    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, 1, DllStructGetSize($tBMI) - 4) ;Structgröße abzüglich der Daten für die Palette
    DllStructSetData($tBMI, 2, $iwidth)
    DllStructSetData($tBMI, 3, -$iheight) ;minus =standard = bottomup
    DllStructSetData($tBMI, 4, 1)
    DllStructSetData($tBMI, 5, 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ückgeben
    EndFunc ;==>_CreateNewBmp32

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

    Func _fps()
    WinSetTitle($hgui, "", $titeltext & $fps & " FPS")
    $fps = 0
    EndFunc

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

    Func _mousedown()
    $mouseflag = 1
    EndFunc

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

    Func _steuerung()
    If WinActive("Apfelmännchen") Then
    If _IsPressed("01", $dll) Then _leftmouse()
    If _IsPressed("02", $dll) Then _rightmouse()
    EndIf
    EndFunc

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

    Func _leftmouse() ;zoom in
    If $mouseflag Then ;neues center
    $posx = MouseGetPos(0)
    $posy = MouseGetPos(1)
    $centerx = ($centerx - $apfelbreite / 2) + $posx * $apfelbreite / $width
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $centerx = ' & $centerx & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $centery = ($centery - $apfelbreite / 2) + $posy * $apfelbreite / $height
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $centery = ' & $centery & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $mouseflag = 0
    EndIf
    $maxiteration /= 0.99785 ;Iterationen erhöhen
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $maxiteration = ' & $maxiteration & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $apfelbreite *= 0.985
    ;~ ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $apfelbreite = ' & $apfelbreite & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    EndFunc

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

    Func _rightmouse() ;zoom out
    If $mouseflag Then ;neues center
    $posx = MouseGetPos(0)
    $posy = MouseGetPos(1)
    $centerx = ($centerx - $apfelbreite / 2) + $posx * $apfelbreite / $width
    $centery = ($centery - $apfelbreite / 2) + $posy * $apfelbreite / $height
    $mouseflag = 0
    EndIf
    $maxiteration *= 0.99785 ;Iterationen erhöhen
    $apfelbreite /= 0.985
    EndFunc

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

    Func _Exit()
    Exit
    EndFunc ;==>_Exit

    [/autoit]

    Kernelintern wird ja width*height berechnet, entspricht also $DATA_SIZE, die Rückgabe des Kernel soll ja nur $DATA_SIZE/2 groß sein - entsprechend müssen die Buffer also auch nur die Größe von $DATA_SIZE/2 haben. Richtig?


    Jein (eigentlich ein klares NEIN 8) )
    Kernelintern wird EIN Pixel berechnet! Anhand der get_global_id(0). Sonst NIX! Punkt!
    WELCHES Pixel berechnet wird, entscheidet ausschliesslich OpenCl. Ansonsten würde niemand ( niemand!!!! ) eine parallelle Berechnung brauchen!
    Ich habe schon OCL-Kernel von gestandenen Programmierern gesehen, die sich darüber ausgelassen haben, dass der Kernel so unglaublich langsam seien. Das lag genau an deinem Problem! Die sind einfach hingegangen und haben den Kernel so geschrieben, dass ALLE Workitems (von ca. 1000 verfügbaren) ALLE Berechnungen für den gesamten Speicherbereich durchgeführt hatten. Das wäre so, als ob jedes einzelne Workitem ALLE Pixel des Apfelmännchens berechnen würde. Und das 1000 mal...umgangssprachlich nennt sich so etwas SCHWACHSINN! Die würden niemals auf die Idee kommen, 1000 Threads anzuwerfen, um in jedem Thread die KOMPLETTE Berechnung durchzuführen. Dieses "Problem" durchzieht die Foren wie ein roter Faden!

    Aber dass jeder einzelne Puffer bei Verwendung von 2 Geräten jeweils nur halb so groß sein muss, hast du ja richtigerweise festgestellt! Wenn man 2x die Hälfte berechnet, kommt immer noch ein Ganzes raus :thumbup:

    Bzgl. deiner Frage zu den Threads, da mache ich mir keinen Kopp! Ob das nun "Parallelisierung" heisst, oder "Threading", ist mir persönlich schnurz. Wichtig ist, dass man begriffen hat, WIE muss mein Kernel aussehen, damit ich das Problem in möglichst VIELE Einzelteile zerlegen kann, welche ALLE gleichzeitig(!) berechnet werden!
    Beispiel: Einzelne Buchstaben in einem String zählen
    Schreib doch mal einen Kernel, der sämtliche "a"´s (oder ggf. nur die Vokale) in einem Text zählt....

    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

    4 Mal editiert, zuletzt von Andy (1. November 2014 um 18:56)

  • Mahlzeit :D

    Ich mache eine neue Meldung, und zwar wollte ich auf meinem System die Intel OpenCL Runtime updaten. Dadurch hat sich die Device Reihenfolge verschoben!
    Und das seltsammste ist jetzt der CL-CPU Aufruf '_CL_GetDevice("cpu", 1,2)'... Wie man entnehmen kann ist die Device und die Platform iwie vertauscht, bei diesem
    Auruf müsste eigentlich der 2te GPU ausgewählt werden. Wenn ich den Aufruf wie gewohnt mit nur '_CL_GetDevice("cpu")' mache, passiert es das automatisch der erste GPU
    ausgewählt wird...
    Ich werde die gesammte OpenCL Installation vom Rechner entfernen und erst CPU und anschliessend GPU installieren, mal sehen ob das dir Reihenfolge wieder herstellt.
    Auch laufen jetzt die von mir Gemoddeten Funktionen nicht mehr! Ebenso deine Multifunktionen.
    Einige Vorteile hat das Update von Intel aber, ich erhalte jetzt etwas mehr Infos in der SciTE Ausgabe und auch ist der CPU bei Simd-Kerneln jetzt etwas schneller (10fps).
    Die Verwendung von 64Bit beschleunigt "_CL_CreateBuffer()" um fast 50%!

    >Device verfügbar= 1;1;4;GPU;GeForce GTX 460;58662400;57218520
    >Device verfügbar= 1;2;4;GPU;GeForce GTX 460;58662480;57218520
    >Device verfügbar= 2;1;2;CPU;Intel(R) Core(TM) i7 CPU 920 @ 2.67GHz;43780328;43771960


    Ich habe auch versucht die angesprochene "Zeichen-Zählfunktion" in OCL zu realisieren, habe aber wieder Probleme mit dem Kernel, bzw. mit der Parameterübergabe. Sieh dir mal bitte den Kernel an, und sag mir was noch immer falsch ist - ich komme einfach nicht drauf?!?!?
    Das ist furchtbar, bei jeder Ausführung erhalte ich untreschiedliche Ergebnisse... Liegt das jetzt am Kernel, oder an der Weise wie ich die Parameter übergebe?

    Spoiler anzeigen
    [autoit]


    #Region - Includes -----------------------------------------------------------------------------------------------------------------
    #AutoIt3Wrapper_UseX64=n
    #AutoIt3Wrapper_UseUPX=n

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

    #RequireAdmin
    #include-once
    #include <opencl_easy.au3>
    #include <GUIConstantsEx.au3>
    #include <WinAPI.au3>

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

    Global _
    $KernelSource, _
    $hTime, $hTimer, $sString, $aData, _
    $sInputFile

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

    #EndRegion -------------------------------------------------------------------------------------------------------------------------

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

    #cs - ReadMe -----------------------------------------------------------------------------------------------------------------------

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

    AutoIt Version: 3.3.12.0
    Script: OCL Test.au3
    Author: Techmix

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

    Programm Funktion:
    OpenCL Testscript "CountChar"

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

    ## --------------------------------------------------------------------------------------------------------------------------------
    ## -Changelog:---------------------------------------------------------------------------------------------------------------------

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

    Funktion1:
    Zählen von Zeichen

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

    #ce --------------------------------------------------------------------------------------------------------------------------------

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

    ; ----------------------------------------------------------------------------------------------------------------------------------
    ; StartUP
    ; ----------------------------------------------------------------------------------------------------------------------------------
    $sInputFile = String(StringToBinary(FileRead(@AutoItExe)))
    dim $aData[] = ["3", "A", "7", "E", "0", "F"]

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

    $hTimer = TimerInit()
    CountChar_OCL($sInputFile, $aData)
    $hTimer = Round(TimerDiff($hTimer),3)
    $sString = "OCL Funktion1 brauchte " & $hTimer & " ms für die Durchführung."
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $sString = ' & $sString & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $hTimer = TimerInit()
    CountChar_AU3($sInputFile, $aData)
    $hTimer = Round(TimerDiff($hTimer),3)
    $sString = "AU3 Funktion1 brauchte " & $hTimer & " ms für die Durchführung."
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $sString = ' & $sString & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console

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

    ;~ Au3 ReplaceFunktion (Vorgabe)
    ;~ 3 94268
    ;~ A 41352
    ;~ 7 79898
    ;~ E 71190
    ;~ 0 487123
    ;~ F 172216

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

    ;~ Ergebnisse:
    ;~ 32 Bit CPU:
    ;~ Kernel Init: 400.655799815618
    ;~ Kernel Time: 1.10142625420035
    ;~ Copy Buffer: 2.92092354687414
    ;~ Kernel Gesammtzeit: 404.700229341138

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

    ;~ 64 Bit CPU:
    ;~ Kernel Init: 189.857790574774
    ;~ Kernel Time: 1.21630482051584
    ;~ Copy Buffer: 2.8847639981453
    ;~ Kernel Gesammtzeit: 193.982539097912

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

    ;~ CPU Sammelergebniss
    ;~ 3 75643
    ;~ A 37321
    ;~ 7 69194
    ;~ E 62256
    ;~ 0 197708
    ;~ F 104693

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

    ;~ 32 Bit GPU:
    ;~ Create Buffer: 407.832135698396
    ;~ Kernel Time: 1.12382273808274
    ;~ Copy Buffer: 2.82555659945063
    ;~ Kernel Gesammtzeit: 411.802314716446

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

    ;~ 64 Bit GPU:
    ;~ Create Buffer: 143.192360565342
    ;~ Kernel Time: 1.13502256605339
    ;~ Copy Buffer: 2.86427600472057
    ;~ Kernel Gesammtzeit: 147.215018777312

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

    ;~ GPU Sammelergebniss
    ;~ 3 1108
    ;~ A 1577
    ;~ 7 1663
    ;~ E 1821
    ;~ 0 2171
    ;~ F 2666

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

    ; ----------------------------------------------------------------------------------------------------------------------------------
    ; Hauptfunktionen
    ; ----------------------------------------------------------------------------------------------------------------------------------
    Func CountChar_OCL($sInput, ByRef $aData)
    ; Zählen von Daten
    $sString = @CRLF & "CountChar_OCL() Zählen von Daten"
    ConsoleWrite($sString & @CRLF)
    Local _
    $aArray[UBound($aData)][2], _
    $sString, _
    $aThreadData, _
    $pThreadStruct, _
    $pOutputSearchStruct, _
    $pOutputCountStruct, _
    $CL_buffer_Input, _
    $CL_buffer_Search, _
    $CL_buffer_Count

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

    for $i = 0 to UBound($aData)-1
    $aArray[$i][0] = $aData[$i]
    $aArray[$i][1] = 0
    Next

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

    $CL_DEBUGFLAG = 0

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

    ; $sInput an Struct übergeben
    $sString = "$sInput in eine Struct übertragen..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    $pThreadStruct = DllStructCreate("char[" & StringLen($sInput) & "]")
    DllStructSetData($pThreadStruct, 1, $sInput)
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    ; $aData an Struct übergeben
    $sString = "$aData in eine Struct übertragen..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    $pOutputSearchStruct = DllStructCreate("char suche[" & UBound($aArray, 1) & "]")
    $pOutputCountStruct = DllStructCreate("int count[" & UBound($aArray, 1) & "]")
    for $i = 0 to UBound($aArray, 1) -1
    DllStructSetData($pOutputSearchStruct, "suche", $aArray[$i][0], $i+1)
    DllStructSetData($pOutputCountStruct, "count", $aArray[$i][1], $i+1)
    Next
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    ; Kernel

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

    ; OK, der Kernel will mich ärgern :D Ich habe diesen so geplant das der "input" aufgeteilt werden soll. Dann sollte alles mit einem Aufruf
    ; abgearbeitet sein. Es sollen also soviele Threads durchlaufen werden, wie es Einträge in "input" gibt, jeder Thread soll jetzt seine "data" mit den
    ; Einträgen in "search" abgleichen, und wenn diese identisch sind soll der entsprechende "count" um einen erhöt werden - dann "return"...
    ; Funzt aber nicht richtig! Die Rückgabe hat immer unterschiedliche Werte?!?
    ; Ich bekomme folgende Meldung "Kernel <counter> was not vectorized" - jetzt befürchte ich das die Threads nicht explizit auf den "input" aufgeteilt
    ; werden, sondern es wird versucht ein Mittelwert aus allen "__global" Parametern zu ermitteln?!?
    ; ...jetzt bekomme im Output schonmal 'Kernel <counter> was successfully vectorized (4)', ist das schonmal gut? :D
    $KernelSource = _
    "__kernel void counter(__global char* input, __constant char* search, __global int* count, const unsigned int iNumElements)" & @CRLF & _
    "{" & @CRLF & _
    " uint threadid = get_global_id(0);" & @CRLF & _
    " char data = input[threadid];" & @CRLF & _
    " uint i;" & @CRLF & _
    " for (i = 0; i < iNumElements; i++) {" & @CRLF & _
    " if (search[i] == data) {" & @CRLF & _
    " count[i] += 1;" & @CRLF & _
    " }" & @CRLF & _
    " }" & @CRLF & _
    "}"

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

    ;~ $KernelSource = _ ; 2.ter Versuch, noch variableres Ergebniss
    ;~ "__kernel void counter(__global char* input, __constant char* search, __global int* count, const unsigned int iNumElements)" & @CRLF & _
    ;~ "{" & @CRLF & _
    ;~ " uint threadid = get_global_id(0);" & @CRLF & _
    ;~ " char data = input[threadid];" & @CRLF & _
    ;~ " uint i, out;" & @CRLF & _
    ;~ " for (i = 0; i < iNumElements; i++) {" & @CRLF & _
    ;~ " if (search[i] == data) {" & @CRLF & _
    ;~ " out = i;" & @CRLF & _
    ;~ " }" & @CRLF & _
    ;~ " }" & @CRLF & _
    ;~ " count[out] += 1;" & @CRLF & _
    ;~ "}"

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

    ;~ $KernelSource = _ ; Simd-Versuch, Crasht rigoros...
    ;~ "__kernel void counter(__global char* input, __global char* search, __global int* count, const unsigned int iNumElements)" & @CRLF & _
    ;~ "{" & @CRLF & _
    ;~ " uint threadid = get_global_id(0)*4;" & @CRLF & _
    ;~ " char data[4], sdata;" & @CRLF & _
    ;~ " data[0] = input[threadid];" & @CRLF & _
    ;~ " data[1] = input[threadid+1];" & @CRLF & _
    ;~ " data[2] = input[threadid+2];" & @CRLF & _
    ;~ " data[3] = input[threadid+3];" & @CRLF & _
    ;~ " uint i, j;" & @CRLF & _
    ;~ " char string;" & @CRLF & _
    ;~ " for (j = 0; j < 4; j++) {" & @CRLF & _
    ;~ " sdata = data[j];" & @CRLF & _
    ;~ " for (i = 0; i < iNumElements; i++) {" & @CRLF & _
    ;~ " string = search[i];" & @CRLF & _
    ;~ " if (string == sdata) {" & @CRLF & _
    ;~ " count[i] += 1;" & @CRLF & _
    ;~ " }" & @CRLF & _
    ;~ " }" & @CRLF & _
    ;~ " }" & @CRLF & _
    ;~ "}"

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

    $hTime = TimerInit()
    _CL_GetDevice("cpu", 1,2)
    ;~ _CL_GetDevice("gpu", 1,1)
    $CL_buffer_Input = _CL_CreateBuffer($pThreadStruct)
    $CL_buffer_Search = _CL_CreateBuffer($pOutputSearchStruct)
    $CL_buffer_Count = _CL_CreateBuffer($pOutputCountStruct)
    _CL_SetArg(0, "ptr*", $CL_buffer_Input)
    _CL_SetArg(1, "ptr*", $CL_buffer_Search)
    _CL_SetArg(2, "ptr*", $CL_buffer_Count)
    _CL_SetArg(3, "int*", UBound($aArray, 1))
    $sString = @CRLF & "Create Buffer: " & TimerDiff($hTime) & @CRLF

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

    $hTime2 = TimerInit()
    _CL_RunKernel(StringLen($sInput), 0)
    $hTime2 = TimerDiff($hTime2)
    $sString &= "Kernel Time: " & $hTime2 & @CRLF
    $hTime2 = TimerInit()
    _CL_ReadBuffer($CL_buffer_Count, $pOutputCountStruct)
    $hTime2 = TimerDiff($hTime2)
    $sString &= "Copy Buffer: " & $hTime2 & @CRLF
    $hTime = TimerDiff($hTime)
    $sString &= "Kernel Gesammtzeit: " & $hTime & @CRLF
    ConsoleWrite($sString & @CRLF)

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

    ; Auswertung
    for $i = 0 to UBound($aArray, 1) -1
    $aArray[$i][1] = DllStructGetData($pOutputCountStruct, 1, $i+1)
    Next
    ;~ _ArrayDisplay($aArray)
    $sString = ""
    for $i = 0 to UBound($aArray, 1) -1
    $sString &= DllStructGetData($pOutputSearchStruct, 1, $i+1) & chr(9) & DllStructGetData($pOutputCountStruct, 1, $i+1) & @CRLF
    Next
    ConsoleWrite($sString & @CRLF)

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

    Return

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

    EndFunc

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

    Func CountChar_OCL2($sInput, ByRef $aData) ; Versuch den Input und Output 'threadid' Syncron zu bekommen, liefert KEIN Ergbniss...
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : StringLeft($sInput,100) = ' & StringLeft($sInput,500) & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    ; Zählen von Daten
    $sString = @CRLF & "CountChar_OCL() Zählen von Daten"
    ConsoleWrite($sString & @CRLF)
    Local _
    $aArray[UBound($aData)], _
    $sString, _
    $aThreadData, _
    $pInputStruct, _
    $pStruct, _
    $pSearchStruct, _
    $pOutputStruct, _
    $CL_buffer_Input, _
    $CL_buffer_Output, _
    $CL_buffer_Search, _
    $CL_buffer_Count

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

    for $i = 0 to UBound($aData)-1
    $aArray[$i]= $aData[$i]
    Next

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

    $CL_DEBUGFLAG = 0

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

    ; $sInput an Struct übergeben
    $sString = "$sInput in eine Struct übertragen..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    $pInputStruct = DllStructCreate("char[" & StringLen($sInput) & "]")
    DllStructSetData($pInputStruct, 1, $sInput)
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    $pOutputStruct = DllStructCreate("int[" & StringLen($sInput) & "]")
    ConsoleWrite($sString & @CRLF)

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

    ; $aData an Struct übergeben
    $sString = "$aData in eine Struct übertragen..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    $pSearchStruct = DllStructCreate("char suche[" & UBound($aArray, 1) & "]")
    for $i = 0 to UBound($aArray, 1) -1
    DllStructSetData($pSearchStruct, "suche", $aArray[$i], $i+1)
    Next
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    ; Kernel

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

    ; OK, der Kernel will mich ärgern :D Ich habe diesen so geplant das der "input" aufgeteilt werden soll. Dann sollte alles mit einem Aufruf
    ; abgearbeitet sein. Es sollen also soviele Threads durchlaufen werden, wie es Einträge in "input" gibt, jeder Thread soll jetzt seine "data" mit den
    ; Einträgen in "search" abgleichen, und wenn diese gleich sind soll der entsprechende "count" um einen erhöt werden - dann "return"...
    ; Funzt aber nicht richtig! Die Rückgabe hat immer unterschiedliche Werte?!?
    ; Ich bekomme folgende Meldung "Kernel <counter> was not vectorized" - jetzt befürchte ich das die Threads nicht explizit auf den "input" aufgeteilt
    ; werden, sondern es wird versucht ein Mittelwert aus allen 3 "__global" Parametern zu ermitteln?!?
    $KernelSource = _
    "__kernel void counter(__global char* input, __global int* output, __constant char* search, const unsigned int iNumElements)" & @CRLF & _
    "{" & @CRLF & _
    " uint threadid = get_global_id(0);" & @CRLF & _
    " char data = input[threadid];" & @CRLF & _
    " uint i, out;" & @CRLF & _
    " for (i = 0; i < iNumElements; i++) {" & @CRLF & _
    " if (search[i] == data) {" & @CRLF & _
    " out = i;" & @CRLF & _
    " }" & @CRLF & _
    " }" & @CRLF & _
    " output[threadid] = out;" & @CRLF & _
    "}"

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

    $hTime = TimerInit()
    _CL_GetDevice("cpu", 1,2)
    ;~ _CL_GetDevice("gpu", 1,1)
    $CL_buffer_Input = _CL_CreateBuffer($pInputStruct)
    $CL_buffer_Output = _CL_CreateBuffer($pOutputStruct)
    $sString = "Create Buffer: " & TimerDiff($hTime) & @CRLF
    _CL_SetArg(0, "ptr*", $CL_buffer_Input)
    _CL_SetArg(1, "ptr*", $CL_buffer_Output)
    _CL_SetArg(2, "ptr*", DllStructGetPtr($pSearchStruct))
    _CL_SetArg(3, "int*", UBound($aArray, 1))
    $sString &= "Kernel Init: " & TimerDiff($hTime) & @CRLF

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

    $hTime2 = TimerInit()
    _CL_RunKernel(StringLen($sInput), 0)
    $hTime2 = TimerDiff($hTime2)
    $sString &= "Kernel Time: " & $hTime2 & @CRLF
    $hTime2 = TimerInit()
    _CL_ReadBuffer($CL_buffer_Output, $pOutputStruct)
    $hTime2 = TimerDiff($hTime2)
    $sString &= "Copy Buffer: " & $hTime2 & @CRLF
    $hTime = TimerDiff($hTime)
    $sString &= "Kernel Gesammtzeit: " & $hTime & @CRLF
    ConsoleWrite($sString & @CRLF)

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

    ; Auswertung
    $sInput = ""
    for $i = 1 to 500
    $sInput &= DllStructGetData($pOutputStruct, 1, $i)
    Next
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : StringLeft($sInput,100) = ' & StringLeft($sInput,500) & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console

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

    EndFunc

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

    Func CountChar_AU3($sInput, ByRef $aData)
    ; Zählen von Daten
    $sString = @CRLF & "CountChar_AU3() Zählen von Daten"
    ConsoleWrite($sString & @CRLF)
    Local _
    $iModeAuto = 1, _ ; $hTime 1391.72808952666, Au3 internal StringReplace()
    $iModeString = 0, _ ; $hTime Unendlich, Zusätzliche Möglichkeiten geschrieben
    $iModeArray = 0, _ ; $hTime 19636.5094663843
    $iModeStruct = 0 ; $hTime 21601.293159999
    Local _
    $aArray[UBound($aData)][2], _
    $sChar, _
    $sString, _
    $aThreadData, _
    $aOutputSearchStruct, _
    $aOutputCountStruct, _
    $pThreadStruct, _
    $pOutputSearchStruct, _
    $pOutputCountStruct

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

    for $i = 0 to UBound($aData)-1
    $aArray[$i][0] = $aData[$i]
    $aArray[$i][1] = 0
    Next

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

    If $iModeAuto = 1 Then
    $sString = "Au3 StringReplace..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    for $i = 0 to UBound($aArray, 1)-1
    StringReplace($sInput, String($aArray[$i][0]), "")
    $aArray[$i][1] = @extended
    Next
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    ElseIf $iModeString = 1 Then
    ; ...Nicht zu empfehlen, braucht ewig!
    $sString = "$iModeString Auswertung..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    for $i = 1 to StringLen($sInput)
    $sChar = StringMid($sInput, $i, 1)
    for $j = 0 to UBound($aArray, 1)-1
    if String($aArray[$j][0]) = $sChar Then
    $aArray[$j][1] += 1
    ContinueLoop 2
    EndIf
    Next
    Next
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    Elseif $iModeArray = 1 then
    ; $sInput an Array übergeben
    $sString = "$sInput an Array übergeben..."
    ConsoleWrite($sString & @CRLF)
    $aThreadData = StringSplit($sInput, "", 2)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    ; Auswertung
    $sString = "$iModeArray Auswertung..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    for $i = 0 to UBound($aThreadData)-1
    $sChar = String($aThreadData[$i])
    for $j = 0 to UBound($aArray, 1)-1
    if String($aArray[$j][0]) = $sChar Then
    $aArray[$j][1] += 1
    ContinueLoop 2
    EndIf
    Next
    Next
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    Elseif $iModeStruct = 1 then
    ; $sInput an Struct übergeben
    $sString = "$sInput in eine Struct übertragen..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    $pThreadStruct = DllStructCreate("char[" & StringLen($sInput) & "]")
    DllStructSetData($pThreadStruct, 1, $sInput)
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    $sString = "$aData in eine Struct übertragen..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    $pOutputSearchStruct = DllStructCreate("char suche[" & UBound($aArray, 1) & "]")
    $pOutputCountStruct = DllStructCreate("int count[" & UBound($aArray, 1) & "]")
    for $i = 0 to UBound($aArray, 1) -1
    DllStructSetData($pOutputSearchStruct, "suche", $aArray[$i][0], $i+1)
    DllStructSetData($pOutputCountStruct, "count", $aArray[$i][1], $i+1)
    Next
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    ; Auswertung
    $sString = "$iModeStruct Auswertung..."
    ConsoleWrite($sString & @CRLF)
    $hTime = TimerInit()
    for $i = 1 to StringLen($sInput)
    $sChar = String(DllStructGetData($pThreadStruct, 1, $i))
    for $j = 0 to UBound($aArray, 1)-1
    if String($aArray[$j][0]) = $sChar Then
    ;~ DllStructSetData($pOutputCountStruct, "count", DllStructGetData($pOutputCountStruct, "count", $j+1)+1, $j+1)
    $aArray[$j][1] += 1
    ContinueLoop 2
    EndIf
    Next
    Next
    $hTime = TimerDiff($hTime)
    $sString = "$hTime " & $hTime
    ConsoleWrite($sString & @CRLF)

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

    EndIf

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

    $sString = ""
    for $i = 0 to UBound($aArray, 1) -1
    $sString &= $aArray[$i][0] & chr(9) & $aArray[$i][1] & @CRLF
    Next
    ConsoleWrite($sString & @CRLF)

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

    EndFunc

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

    Exit

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

    Grüsse!

    P.S: Entschuldigt die schreckliche Formatierung, wurde mit dem Handy geschrieben...

    EDIT
    Ich vergaß die SciTE Consolenausgabe...

    Spoiler anzeigen

    Einmal editiert, zuletzt von Techmix (5. November 2014 um 15:35)

  • >Device verfügbar= 1;1;4;GPU;GeForce GTX 460;58662400;57218520
    >Device verfügbar= 1;2;4;GPU;GeForce GTX 460;58662480;57218520
    >Device verfügbar= 2;1;2;CPU;Intel(R) Core(TM) i7 CPU 920 @ 2.67GHz;43780328;43771960

    Ist das die Ausgabe vor oder nach der Neuinstallation? Ist so doch einwandfrei!

    Übrigens hatte ich zufällig letzte Woche einen Rechner zum Neuaufsetzen da, und weil der auch mit Intel-CPU und Nvidia-Graka daherkam, habe ich OpenCl gleich ausprobiert.
    Treiber installiert, läuft.
    Sowohl die "normalen" Beispiele, als auch die Multifunktionen.
    Du musst natürlich aufpassen, dass du die richtigen command_queue[Platform][Device]-Indizes setzt! Wenn vorher 2 Devices(GPUs) auf der "ersten" Platform gesessen haben, und nach dem Installieren dann die CPU "solo", dann stimmen natürlich die Parameter nicht!
    Das hat aber weniger mit OpenCl zu tun, sondern mit deiner Implementation.

    Natürlich könnte man auf dem Standpunkt stehen, eine "Engine" sollte die Last auf alle Devices selbst verteilen können....
    Wie schon weiter oben gesagt, ist Lastverteilung die Sache des Anwenders.
    "Professionelle" Programme machen das übrigens so, dass bei SLI-Verbünden einfach in einen komplett anderen Programmzweig gesprungen wird. D.h. dort gehen die Programmierer den pragmatischen Weg und testen auf Platform1/CPU_oder_SLI und Platform2/CPU_oder_SLI. Problem gelöst^^
    Kannst du übrigens auch so machen, musst nur einen Zweig einbauen für
    _CL_RunKernel_ALL($DATA_SIZE / 2, 0, $comm_queue[1][2]) bzw
    _CL_RunKernel_ALL($DATA_SIZE / 2, 0, $comm_queue[2][2]) je nachdem, ob die GPUs auf Platform 1 oder 2 sitzen

    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

    Einmal editiert, zuletzt von Andy (6. November 2014 um 18:00)

  • Sodele, habe mal deinen Kernel laufen lassen, ich hoffe, dir gehen jetzt Kronleuchter auf^^

    &quot;Script&quot;
    [autoit]

    #AutoIt3Wrapper_UseX64=n

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

    #include <opencl_easy.au3>

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

    $KernelSource = _
    "__kernel void counter(__global char* input, __constant char* search, __global int* count, const unsigned int iNumElements)" & @CRLF & _
    "{" & @CRLF & _
    " uint threadid = get_global_id(0);" & @CRLF & _
    " char data = input[threadid];" & @CRLF & _
    " uint i;" & @CRLF & _
    " for (i = 0; i < iNumElements; i++) {" & @CRLF & _
    " if (search[i] == data) {" & @CRLF & _
    " count[i] += 1;" & @CRLF & _
    " }" & @CRLF & _
    " }" & @CRLF & _
    "}"

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

    $sInputFile = String(StringToBinary(FileRead(@AutoItExe)))
    $DATA_SIZE = StringLen($sInputFile) ;bildgröße=anzahl pixel (rgba)
    $aData = "3A7E0F"
    $Elements = StringLen($adata)

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

    _CL_GetDevice("GPU", 1,1) ;..oder CPU oder GPU, ermittelt das Gerät, auf dem die Berechnungen durchgeführt werden.

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

    ;Puffer für das file
    $buffer_data = _dllstructcreate16("char[" & $DATA_SIZE & "]");16byte-aligned!
    $CL_buffer_data = _CL_CreateBuffer($buffer_data)
    ;Puffer für den Suchstring
    $buffer_search = _dllstructcreate16("char[" & $elements & "]")
    $CL_buffer_search = _CL_CreateBuffer($buffer_search)
    ;Puffer für die Rückgabe
    $Buffer_count = _dllstructcreate16("int[" & $elements & "]")
    $CL_buffer_count = _CL_CreateBuffer($Buffer_count)

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

    ;Puffer füllen
    DllStructSetData($buffer_data, 1, $sInputFile)
    _CL_WriteBuffer($CL_buffer_data, $buffer_data)

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

    DllStructSetData($buffer_search, 1, $aData)
    _CL_WriteBuffer($CL_buffer_search, $buffer_search)

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

    ;Parameter an den Kernel übergeben
    _CL_SetArg(0, "ptr*", $CL_buffer_data)
    _CL_SetArg(1, "ptr*", $CL_buffer_search)
    _CL_SetArg(2, "int*", $CL_buffer_count)
    _CL_SetArg(3, "int*", $Elements)

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

    $CL_DEBUGFLAG = 1 ;Debug-Ausgabe mit 0 ausschalten, um Geschwindigkeit zu erhöhen
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $CL_DEBUGFLAG = ' & $CL_DEBUGFLAG & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console

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

    $t = TimerInit()
    _CL_RunKernel($DATA_SIZE, 0)
    $m = TimerDiff($t)
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $m = ' & $m & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    _CL_ReadBuffer($CL_buffer_count, $Buffer_count)
    $m = TimerDiff($t)
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $m = ' & $m & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console

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

    For $i = 1 To $Elements
    $buchstabe=stringmid($aData,$i,1)
    stringreplace($sInputFile,$Buchstabe,$Buchstabe,0,1) ;AutoIt rulez^^
    $ext = @extended
    $ret = DllStructGetData($Buffer_count, 1, $i)
    ConsoleWrite($Buchstabe & " OCl=" & $ret &" AutoIt="&$ext & @CRLF) ;### Debug Console
    Next

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

    ;...das wars schon^^

    [/autoit]
    &quot;Ergebnisse&quot;
    [autoit]

    verwendetes Device 1;1;4;GPU;BeaverCreek;45838144;1558332692

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

    clEnqueueNDRangeKernel CL_SUCCESS
    @@ Debug(55) : $m = 0.753986456087799
    clEnqueueReadBuffer CL_SUCCESS
    @@ Debug(58) : $m = 28.5525658254356
    3 OCl=7886 AutoIt=95934
    A OCl=7100 AutoIt=40589
    7 OCl=9042 AutoIt=80065
    E OCl=8995 AutoIt=74060
    0 OCl=10088 AutoIt=485313
    F OCl=9419 AutoIt=174165

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

    verwendetes Device 1;2;4;GPU;Caicos;58117512;1585136916

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

    clEnqueueNDRangeKernel CL_SUCCESS
    @@ Debug(55) : $m = 1.01703906807567
    clEnqueueReadBuffer CL_SUCCESS
    @@ Debug(58) : $m = 29.2156928835109
    3 OCl=12268 AutoIt=95934
    A OCl=10193 AutoIt=40589
    7 OCl=13415 AutoIt=80065
    E OCl=13260 AutoIt=74060
    0 OCl=15621 AutoIt=485313
    F OCl=13984 AutoIt=174165

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

    verwendetes Device 1;3;2;CPU;AMD A6-3400M APU with Radeon(tm) HD Graphics;72411976;1558332692

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

    clEnqueueNDRangeKernel CL_SUCCESS
    @@ Debug(55) : $m = 0.69097106714363
    clEnqueueReadBuffer CL_SUCCESS
    @@ Debug(58) : $m = 78.0753341647457
    3 OCl=95424 AutoIt=95934
    A OCl=40521 AutoIt=40589
    7 OCl=79568 AutoIt=80065
    E OCl=73754 AutoIt=74060
    0 OCl=396277 AutoIt=485313
    F OCl=169821 AutoIt=174165

    [/autoit]

    Dein Problem mit den "falschen" Zählergebnissen war natürlich schon bei der Aufgabenstellung klar! Sorry, aber manchmal muss man mit der Brechstange vorgehen 8)
    Dein Kernel macht doch genau, was er soll, es wird ein einzelnes Zeichen aus dem großen String ausgelesen, geschaut, ob es im Suchstring enthalten ist, und dann der Zähler für dieses Zeichen erhöht...

    Ich fange bei der Analyse der Ergebnisse mit der CPU an, die kommen der "richtigen" Ergebnissen schon sehr nahe!
    Stell dir mal eine Schneeballschlacht vor. Du bist eine einzelne Speicherstelle, die bei einem Treffer eine Strichliste führen soll. Dir gegenüber stehen 2 (Dualcore) oder 4(Quadkore) "Kerne", welche dich mit Schneebällen bewerfen. Die Strichliste ist relativ einfach zu führen, denn die 4 "Kerne" müssen ja den Schneeball formen und werfen. Ab und zu sind die Jungs aber geringfügig schneller oder langsamer und mehrere Schneebälle treffen dich GLEICHZEITIG! Da ist es natürlich schwierig für dich die genaue Anzahl der gleichzeitigen Treffer zu zählen, und du machst nur EINEN Strich! Das kommt allerdings nicht oft vor...

    Jetzt spielst du gegen eine GPU mit hunderten, wenn nicht sogar tausenden "Kernen". Alle werfen GLEICHZEITIG mit Schneebällen nach dir und du sollst JEDEN Treffer zählen....unmöglich?! Genau! Das kann niemand, auch die Speicherstelle im Rechner nicht! Capice?!

    DAS ist paralleles Computing!

    Um die "richtige" Trefferanzahl herauszubekommen, musst du jetzt nur noch herausfinden, wie man die "Schneeballwerfer" synchronisiert :D Schönes Wochenende! :thumbup:

  • Lösung sind die sog. "atomic operations".
    Hier und hier gibts dazu bissl was zu lesen...

    Problem dabei ist, dass sämtliche "atomics" sehr (sehr) zeitaufwendig, und damit langsam ausgeführt werden!
    Um beim Beispiel mit der Schneeballschlacht zu bleiben, wenn eine Million Leute nacheinander werfen müssen, dann dauert das WESENTLICH länger, als wenn alle gleichzeitig werfen!
    Atomic operations sind also das, was man bei parallelem Computing unbedingt vermeiden soll!

    Um es nochmal deutlich zu sagen, OpenCL ist das Werkzeug für massives paralleles Number crunching. Um bissl pillepalle Zählaufgaben wie im Beispiel abzuwickeln, braucht man das wirklich nicht!
    Allerdings ist das ein sehr schönes Beispiel, wie man lokalen Speicher, barriers, und auch atomic operations einsetzen kann. Und die Aufteilung in Teilstücke ^^

    Um die Verwendung von atomics dem Kernel mitzuteilen, benutzt man die #pragma-Anweisung. Genau wie die Verwendung von printf() zur bspw. Debug-Ausgabe direkt aus dem Kernel in die Console! printf() sollte man, sobald der Kernel fertig entwickelt ist, tunlichst wieder entfernen!

    Anbei eine Methode, wie ich für MEINE Devices die optimalen Teilgrößen berechne. Alternativ kann man natürlich auch problemlos beim _CL_RunKernel$DATA_SIZE, 0) statt der Null (OCL legt selbstständig die optimale Workgroupsize fest) die Workgroupsize festlegen.
    Mit den somit herausgefundenen Parametern hat man dann für SEINE Hardware die schnellste Software am Start...

    Spoiler anzeigen
    [autoit]

    #AutoIt3Wrapper_UseX64=n ;

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

    #include <opencl_easy_all.au3>

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

    $KernelSource = _
    "#pragma OPENCL EXTENSION cl_amd_printf : enable " & @CRLF & _;gibts auch für INTEL/NVIDIA, zum debuggen
    "#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable " & @CRLF & _;atomics einschalten
    "__kernel void counter(__global char* input, __constant char* search, __global uint* count, const int iNumElements,const int divisor)" & @CRLF & _
    "{ " & @CRLF & _
    " uint threadid = get_global_id(0)*divisor; " & @CRLF & _;später "Pointer" auf den Anfang der Teilstrings
    " uint i,r; " & @CRLF & _;For/To variablen
    " private uint local_array[16] ; " & @CRLF & _;LOKALES Array , dann muss nicht immer aus dem Globalen Speicher gelesen werden!
    " for (i = 0; i < iNumElements; i++) { " & @CRLF & _;Lokales array nullen....Lokaler Speicher wird nicht initialisiert!
    " local_array[i]=0; " & @CRLF & _
    " } " & @CRLF & _
    " for (r = 0; r < divisor; r++) { " & @CRLF & _;Teilstrings einlesen
    " char data = input[threadid+r]; " & @CRLF & _;einzelnes Zeichen
    " for (i = 0; i < iNumElements; i++) { " & @CRLF & _;ist es Teil des Suchstrings?
    " if (search[i] == data) { " & @CRLF & _
    " local_array[i]+=1; " & @CRLF & _;ja, dann Anzahl im LOKALEN Array erhöhen
    ' //printf("threadid=%s i=%u loc_arr=%u \n",search[i]);//,i,local_array[i]);' & @CRLF & _;nur zum debuggen
    " } " & @CRLF & _
    " } " & @CRLF & _
    " } " & @CRLF & _
    "barrier(CLK_GLOBAL_MEM_FENCE); " & @CRLF & _;warten, bis alle Workitems fertig sind!
    " for (i = 0; i < iNumElements; i++) { " & @CRLF & _;dann Ergenisse vom LOKALEN Speicher in den Globalen schreiben
    ' //printf(" count[i]=%i i=%u loc_arr=%u \n",count[i],i,local_array[i]);' & @CRLF & _
    " atomic_add(&count[i],local_array[i]); // += 1; " & @CRLF & _;
    " } " & @CRLF & _
    "}"

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

    $sInputFile = String(StringToBinary(FileRead(@AutoItExe))) ;Datei einlesen, Größe als Vielfaches von 2^16
    $DATA_SIZE = StringLen($sInputFile)
    $DATA_SIZE = Int($DATA_SIZE / 2 ^ 20) * 2 ^ 20 ;vielfaches von 2, schöner zu teilen^^

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

    $sInputFile = StringLeft($sInputFile, $DATA_SIZE) ;vielfaches von 2
    $aData = "3A7E0F" ;zu suchende Buchstaben testen! "01234567890ABCDEF"
    $Elements = StringLen($aData) ;Anzahl zu suchende Buchstaben

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

    _CL_GetDevice("ALL", 1, 1) ;..oder CPU oder GPU, 1.Var Devicenummer, 2.Var Platform

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

    ;Puffer für das file
    $buffer_data = _dllstructcreate16("char[" & $DATA_SIZE & "]") ;16byte-aligned!
    $CL_buffer_data = _CL_CreateBuffer($buffer_data)
    ;Puffer für den Suchstring
    $buffer_search = _dllstructcreate16("char[" & $Elements & "]")
    $CL_buffer_search = _CL_CreateBuffer($buffer_search)
    ;Puffer für die Rückgabe
    $Buffer_count = _dllstructcreate16("uint[" & $Elements & "]")
    $CL_buffer_count = _CL_CreateBuffer($Buffer_count)

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

    ;Puffer füllen
    DllStructSetData($buffer_data, 1, $sInputFile)
    _CL_WriteBuffer($CL_buffer_data, $buffer_data)

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

    DllStructSetData($buffer_search, 1, $aData)
    _CL_WriteBuffer($CL_buffer_search, $buffer_search)

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

    ;Parameter an den Kernel übergeben
    _CL_SetArg(0, "ptr*", $CL_buffer_data)
    _CL_SetArg(1, "ptr*", $CL_buffer_search)
    _CL_SetArg(2, "ptr*", $CL_buffer_count)
    _CL_SetArg(3, "int*", $Elements)

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

    $min_time = 2 ^ 31 ;kleinste benötigte Kerneltime
    $min_divisor = 0

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

    $CL_DEBUGFLAG = 0 ;Debug-Ausgabe mit 0 ausschalten, um Geschwindigkeit zu erhöhen

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

    For $p = 0 To 16 ;Divisor ermitteln

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

    $divisor = 2 ^ $p
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $divisor = ' & $divisor & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console

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

    For $w = 1 To $Elements ;zählpuffer nullen
    DllStructSetData($Buffer_count, 1, 0, $w)
    Next
    _CL_WriteBuffer($CL_buffer_count, $Buffer_count) ;und Puffer mit Nullen füllen

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

    _CL_SetArg(4, "int*", $divisor) ;Divisor an Kernel übergeben

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

    _CL_RunKernel(Int($DATA_SIZE / $divisor), 0)
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $DATA_SIZE / $divisor = ' & $DATA_SIZE / $divisor & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console
    $r = profile_cl_event()
    ConsoleWrite('@@ Debug(' & @ScriptLineNumber & ') : $r = ' & $r & @CRLF & '>Error code: ' & @error & @CRLF) ;### Debug Console

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

    If $r < $min_time Then ;besten Divisor finden
    $min_time = $r
    $min_divisor = $divisor
    MsgBox(262144, Int($min_time) & " ms", "Divisor = " & $min_divisor, 2) ;### Debug MSGBOX
    EndIf

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

    _CL_ReadBuffer($CL_buffer_count, $Buffer_count) ;Puffer auslesen

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

    For $i = 1 To $Elements ;und Einzelwerte extrahieren
    $buchstabe = StringMid($aData, $i, 1)
    StringReplace($sInputFile, $buchstabe, $buchstabe, 0, 1) ;AutoIt rulez^^
    $ext = @extended ;anzahl durch AutoIt ersetzter Buchstaben
    $ret = DllStructGetData($Buffer_count, 1, $i) ;Puffer auslesen
    ConsoleWrite($buchstabe & " OCl=" & $ret & " AutoIt=" & $ext & @CRLF) ;und anzeigen
    Next

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

    Next
    ;...das wars schon^^

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

    MsgBox(0, "Testende", "Schnellster Divisor = " & $min_divisor & @CRLF & "mit Teilstringgröße = " & ($DATA_SIZE / $min_divisor) & " dauert " & Int($min_time) & "ms")

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

    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

    Einmal editiert, zuletzt von Andy (10. November 2014 um 07:14)

  • Eine weitere Option wär der Reduce-Ansatz.
    Kann man prinzipiell mit allen assoziativen Operationen durchführen (Hier ist es die Addition).
    Dabei fängt man z.B. mit zwei benachbarten Threads an. Diese bilden gemeinsam eine Summe. Sie müssen also nur auf ihre beiden Threads warten - alle anderen schreiben ja auf eine andere Stelle und machen das parallel zu unseren ersten beiden.
    Dann wird die Summe der ersten beiden mit den von Thread 3-4 gebildet. Dann diese wieder mit Thread 5-8 usw.
    Bis am Ende nur ein Wert übrig bleibt.
    So kann man die Summe partiell parallel berechnen lassen und auf atomare Operationen verzichten.
    Man ist dann weiterhin schneller als die Summe seriell zu berechnen.

    Hier mal zur Verdeutlichung: >>http://web.engr.oregonstate.edu/~mjb/cs575/Handouts/opencl.reduction.2pp.pdf<<

  • Ja, das geht natürlich auch, statt der atomics hat man dann die barriers...und muss zusätzlich nach dem Lesen des Output-Speicherbereichs noch alle Summen zusammenzählen. Wenn die "eigentliche" Anwendung nicht massiv parallelisierbare und bestenfalls float-Arithmetik nutzt, hilft das sicher auch nicht viel.

    Alles in allem stellt sich die Frage, ob man für solch triviale Aufgaben nicht eine CPU und SIMD mit einer SoA-Datenstruktur benutzt.
    Je nach Anwendung rockt dann ein CPU-Kern die tausenden GPU-Kerne. Cést la vie.

  • Ist schon klar :D
    Übrigens, u.a. lesenswert zum Thema SoA/AoS ein INTEL-Paper, welches eine 16-Core XEON-CPU mit einem 61-Core XEON-PHI vergleicht.
    Die Essenz "....if a code yields no vector speedup (no auto-vectorization that produces speedup) at
    all, then it is not very likely that Xeon Phi will ever exceed the performance of Intel Xeon even if it
    is fully parallel. That is, “highly parallel” is not enough to allow Phi to provide benefit in most cases.
    "
    lässt sich imho gut auf die Problematik "Code auf CPU oder GPU?" anwenden.

    Voraussetzung für gut auf der GPU verarbeitbaren Code ist Vektorisierung (4/8/16 "Schneebälle" werden durch SIMD/AVX IMMER gleichzeitig geworfen) , und zusätzlich noch Parallelisierbarkeit über den Algorithmus.


    @Techmix, ich glaube, nun ist es an der Zeit deine Idee von einem Packer/Komprimierer umzusetzen 8o