OpenCl goes AutoIt Update 31.Dezember 2016

Techmix

Hallo!

Danke für die Hinweise zu dem Kernel und der Paralellisierung!! Hatte mich schon gefragt wieso jedes mal andere Ergebnisse zustande kommen - Der Kernel selber konnte es nicht sein, das gibt es nichts zum falsch machen... Aber es wird jetzt immer besser mit OCL

>Gibst mir so eine Aufgabe, und wartest nichtmals das Wochenende ab...< Hatte mich extra damit beschäftigt weil ich mich darüber geärgert hatte, Habe 3 weitere Kernel geschrieben und den 4ten spare ich mir jetzt... Ich hänge das Script jetzt aus Trotz an
Ausserdem habe ich den 2ten Kernel vom vorigen Script lauffähig bekommen Es wird endlich richtig gezählt, aber die Auswertung folgt noch in Au3 und benötigt entsprechend viel Zeit. Gibt es vielleicht noch einen besseren (schnelleren) Weg die Daten zu sortieren?
Ich kläre diese Woche das Kompressionsscript, und prüfe mal welche Teile sich dafür eignen. Das war ja eh in meinem Sinn das Ding Aufzurollen. Die Standartmethode sollte jedoch gut umsetzbar sein, aber ich fände es trotzdem schön so nah an dem Orginal wie möglich zu bleiben...

Wenn man eine Lastanalyse für einen Kernel auf allen Devices durchführen möchte, dann muß ich doch so vorgehen das der Kernel auf jedem Device einzeln für nur einige Durchläufe ausgeführt werden soll - dabei die Zeit stoppen. Kann ich das z.B. so in etwa machen '_CL_RunKernel(StringLen($sInput) / 100, 0)' oder muß ich den Kernel extra dafür anpassen?
Theoretisch kann man Kernel im vorraus so Planen, das diese Multi fähig sind. Aber ich würde noch gerne 'Clustern' also alle OCL-Devices nutzen können (CPU/GPU Mix) - ich weiß inzwischen auch wie ich mit diesen umgehen muß

OpenCL Update

Hatte mich vielleicht undeutlich ausgedrückt
Vor dem Update:

Code

>Device verfügbar= 1;1;2;CPU;Intel(R) Core(TM) i7 CPU         920  @ 2.67GHz;18910080;18975064
>Device verfügbar= 2;1;4;GPU;GeForce GTX 460;57417256;57417200
>Device verfügbar= 2;2;4;GPU;GeForce GTX 460;9218672;57417200
CPU: _CL_GetDevice("cpu", 1,1)
GPU1: _CL_GetDevice("gpu", 2,1)
GPU2: _CL_GetDevice("gpu", 2,2)

Nach Update - Versuch 1:

Code

>Device verfügbar= 1;1;4;GPU;GeForce GTX 460;57417256;57417200
>Device verfügbar= 1;2;4;GPU;GeForce GTX 460;9218672;57417200
>Device verfügbar= 2;1;2;CPU;Intel(R) Core(TM) i7 CPU         920  @ 2.67GHz;18910080;18975064
CPU: _CL_GetDevice("cpu", 1,2)
GPU1: _CL_GetDevice("gpu", 1,1)
GPU2: _CL_GetDevice("gpu", 2,1)

Nach Update - Versuch 2:
Update auf: Nvidia (344.60), Intel (intel_sdk_for_ocl_applications_2014_x64):

Code

32Bit:
>Device verfügbar= 1;1;2;CPU;Intel(R) Core(TM) i7 CPU         920  @ 2.67GHz;46960640;46867144
>Device verfügbar= 2;1;4;GPU;GeForce GTX 460;69475744;69468584
>Device verfügbar= 2;2;4;GPU;GeForce GTX 460;69475800;69468584


64Bit:
WARNING: Linking two modules of different data layouts!
WARNING: Linking two modules of different target triples: C:\Program Files (x86)\Common Files\Intel\OpenCL\bin\common\clbltfnshared.rtl: 'i686-pc-win32-elf' and 'x86_64-pc-win32-elf'
>Device verfügbar= 1;1;2;CPU;Intel(R) Core(TM) i7 CPU         920  @ 2.67GHz;43352304;42695744
>Device verfügbar= 2;1;4;GPU;GeForce GTX 460;45596064;44965392
>Device verfügbar= 2;2;4;GPU;GeForce GTX 460;45596176;44965392


CPU: _CL_GetDevice("cpu", 1,1)
GPU1: _CL_GetDevice("gpu", 1,2)
GPU2: _CL_GetDevice("gpu", 2,2)

Alles anzeigen

Nach erneuter Neu-Installation aller OpenCL Komponenten, hat sich der Fehler immer noch nicht ganz geklärt...
So wie ich das sehe, hat das Update die komplette OCL-Installation durcheinander gebracht...
$num_device_used und $num_platform_used sind vertauscht, deswegen geht kein _cl_GetALLDevice("multi") mehr bei mir!
Seltsammerweise funktioniert aber _CL_GetALLDevice2("multi") mit $all=1 bei mir, _CL_GetALLDevice2("cpu") und $all=0 jedoch nicht... (aus 'tunnel_forum_ALL2.au3')
Die Aufrufe für die Grafikkarten muss ich umbedingt wieder so hinbiegen, wie sie zum Anfang waren - Das sorgt für Probleme bei der Multi Verwendung! Kennst du vielleicht irgendwelche Reg-Key´s (oder so) wo ich händisch eintragen kann wie die DeviceID´s auf die Platformen verteilt werden???

Grüsse!

P.S:
Diesmal hänge ich das Script an, zerstört sonst immer die Formatierung... Die Funktion 'CountChar_OCL2()' enthällt den zählenden Kernel, die anderen Funktionen sind experimentell.

**Andy**

Zitat von Techmix

>Gibst mir so eine Aufgabe, und wartest nichtmals das Wochenende ab...

Der Standardspruch einer meiner Profs vor 30 Jahren zum Theme "zu wenig Zeit" war: "Was macht ihr nachts?"

Zitat von Techmix

Wenn man eine Lastanalyse für einen Kernel auf allen Devices durchführen möchte, dann muß ich doch so vorgehen das der Kernel auf jedem Device einzeln für nur einige Durchläufe ausgeführt werden soll - dabei die Zeit stoppen.

Ja, zwei Möglichkeiten: Erstens die Datengröße in "Teilstücke" ändern, dazu muss dem Kernel über einen Parameter der "Teiler" mitgegeben werden, problemlos!
Zweitens die Anzahl der Workgroups über den zweiten Parameter bei _CL_RunKernel() (ist Standardmäßig null) einstellen. Das ist aber "finetuning", wenn OCL die verwendeten Workgroups festlegt, ist man eigentlich immer nah am Optimum!

Zitat von Techmix

Theoretisch kann man Kernel im vorraus so Planen, das diese Multi fähig sind.

Das wäre natürlich optimal, macht aber nur dann Sinn, wenn die Kernellaufzeiten (reine Berechnungen auf der GPU) so lange dauern, dass der Overhead durch das weitere Übertragen des Speichers , also _CL_ReadBuffer(), nicht den Gewinn wieder auffrisst! Macht imho nur bei wirklich aufwendigen Berechnungen Sinn!

Techmix

Nabend

Ich war ansatzweise Fleißig und bin vier Baustellen gleichzeitig angegangen... Erstmal habe ich für den Kompressionsalgorythmus die absolute Mindest-Basis geschaffen, diese enthällt eine funktionierende Vorlage in AU3 und eine Interpretation in OpenCL. Dann habe ich die Herkunft des Alogo´s entstaubt und einigermaßen Lauffähig bekommen dazu einen Post verfasst und werde diesen noch aktualisieren. Auch habe ich ein Syntax-Highlighting für OpenCL in SciTE erstellt, damit mir das Arbeiten damit leichter fällt - ist aber farblich nicht so dolle geworden...

Meine Motivation liegt darin eine OpenCL Interpretation der Kompressionsmethoden aus cHEX zu realisieren, weil ich die Geschwindigkeitsvorteile durchaus zu schätzen weiß! Außerdem bietet sich das Projekt für OCL an, weil ich dort mit Multiprocessing arbeite; in OCL könnten diese jeweils ein Device darstellen, und die Geschwindigkeit maximieren OHNE dabei auf Multi-Kernel zugreifen zu müssen (Die Verwaltung übernimmt das Main-Script).

Ich habe eine lauffähige Kurzversion, welche das Multiprocessing verwendet, für das Anwendungsbeispiel angefertigt... So bin ich in der Lage innerhalb von AU3 die maximale Threadanzahl zu nutzen und in OpenCL habe ich den Vorteil alle vorhanden Devices als 'Threads' verwalten zu können <- Durch die Multiprocessing-Methode von cHEX wird der Prozess einmalig kompiliert und mehrfach ausgeführt, dabei wird dieser über TCP mit allen nötigen Parametern versorgt, und das Main-Script übernimmt dabei die Prozess- und Clusterverwaltung. Dann kann ein identischer Kernel auf allen Devices mit den jeweiligen Parametern ausgeführt werden - ohne das ich an den Kernel Anpassungen durchführen muss! (Hoffentlich :D)

Ich hänge alles an, aber Poste den Kernel noch - vielleicht kann mir ja jemand ja so schon helfen

Kernel

Code

/*
    SimpleCompression:


    Master:
    Im Master-Script werden die Eingangsdaten in Daten-Cluster aufgeteilt (Momentan 25KB, muß noch optimiert werden)
    Die Daten-Cluster werden nacheinander an den OpenCL-Kernel übergegeben


    Kernel-StartUP:
    Der Kernel teilt seine Eingangsdaten in '$iGlobalSize' große Teile in 'Local_String' auf - Umwandlung dabei nach uInt (wie bei Asc())
    Für alle wichtigen Arrays werden 'Lokale' erstellt und rein kopiert, so soll


    Analyse:
    Char Wiederholungen in 'Local_iNewAnalyse' zählen
    Analyse von jeweils 2 'Local_String's, wenn nicht in 'NotReplace' vorhanden 'Local_iAnalyse[i]+=1'
    ? Analyse-Ergebniss in globalen Speicher addieren
    Fertiggestelltes Analyse-Ergebniss in 'Local_iNewAnalyse' zurück kopieren
    ? 'Local_iNewAnalyse' absteigend sortieren


    Replacing:
    ? 'Local_iNewAnalyse' tauschen mit (uInt)Chars welche nicht in 'Local_NotReplace' vorhanden sind


    Rückgabe:
    ???


*/


#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable    // Atomics einschalten


__kernel void SimpleCompression(
                __constant char* _In, __global char* _Out, const uint iNumElements,            // In- und Output
                __constant uint* NotReplace, const uint iNumNotReplace,                        // NotReplace
                __constant char* Analyse, __global uint* iAnalyse)                            // Analyse
{
// Deklarationen
    uint iGlobalSize    = $iGlobalSize;                                // Eingangs-Datengröße für Globale Daten ($iGlobalSize wird in AU3 Replaced)
    uint iLocalSize        = $iLocalSize;                                // Datengröße je Thread für Lokale Daten ($iLocalSize wird in AU3 Replaced)
    uint iNumAnalyse    = $iNumAnalyse;                                // Anzahl von Analyse-Einträgen ($iNumAnalyse wird in AU3 Replaced)
    uint iThreadID        = get_global_id(0)*iGlobalSize;                // Global-ThreadID alle iDataSize-Char
    uint iLocalID        = get_local_id(0)*2;                        // Local-ThreadID alle 2-Char
    uint iOutput        = 0;                                        // Output Counter
    uint i, j;


// Lokale Array´s
    private uint Local_String[$iGlobalSize] ;                        // Lokales Array 'Local_String'        ($iGlobalSize wird in AU3 Replaced)
    private uint Local_NotReplace[32] ;                                // Lokales Array 'Local_NotReplace'
    private uint Local_sAnalyse[$iNumAnalyse*2] ;                    // Lokales Array 'Local_sAnalyse'    ($iNumAnalyse wird in AU3 Replaced)
    private uint Local_iAnalyse[$iNumAnalyse] ;                        // Lokales Array 'Local_iAnalyse'
    private uint Local_iNewAnalyse[$iNumAnalyse][2] ;                // Lokales Array 'Local_iNewAnalyse'


// Debugging
    //~ for (i = 0; i < iNumAnalyse; i++) {
        //~ printf("i=%i, Analyse[i]=%c \n", i, Analyse[i]);
    //~ }
    for (i = 0; i < iNumAnalyse; i++) {                                // Hier sollte die Ausgabe immer 0 sein...
        printf("i=%i, iAnalyse[i]=%i \n", i, iAnalyse[i]);
    }


    //~ printf("get_global_id(0)=%i \n", get_global_id(0));
    //~ printf("get_global_size(0)=%i \n", get_global_size(0));
    //~ printf("get_local_id(0)=%i \n", get_local_id(0));
    //~ printf("get_local_size(0)=%i \n", get_local_size(0));
    //~ printf("iGlobalSize=%i \n", iGlobalSize);
    //~ printf("iNumElements=%i \n", iNumElements);
    //~ printf("iNumNotReplace=%i \n", iNumNotReplace);
    //~ printf("iNumAnalyse=%i \n", iNumAnalyse);


//
//     Kernel-StartUP


//     Input-Daten verarbeiten...
    for (i = 0; i < iGlobalSize; i++) {                                // 'iGlobalSize' Chars von 'String' in lokales Array kopieren
        if (iThreadID+i < iNumElements) {                            // Aufpassen das 'iThreadID+i' < 'iNumElements' bleibt
            Local_String[i] = convert_int(_In[iThreadID+i]);        // Umwandlung von Char zu Int
            //~ printf("iThreadID+i=%i, _In[iThreadID+i][i]=%c, Local_String[i]=%i \n", iThreadID+i, _In[iThreadID+i], Local_String[i]);
            //~ printf("_In[iThreadID+i]=%c, Local_String[i]=%i, Local_String[i]=%c \n", _In[iThreadID+i], Local_String[i], Local_String[i]);
            //~ printf("_In[iThreadID+i]=%c \n", _In[iThreadID+i]);
        }
    }
    for (i = 0; i < iNumNotReplace; i++) {                            // NotReplace in lokales Array kopieren
        Local_NotReplace[i] = NotReplace[i];
        //~ // printf("Local_NotReplace[i]=%i \n", Local_NotReplace[i]);
    }
    for (i = 0; i < iNumAnalyse*2; i++) {                            // Analyse in lokales Array kopieren
        Local_sAnalyse[i] = convert_int(Analyse[i]);
        //~ // printf("Local_sAnalyse[i]=%i, Local_sAnalyse[i]=%c \n", Local_sAnalyse[i], Local_sAnalyse[i]);
    }
    for (i = 0; i < iNumAnalyse; i++) {                                // Analyse-Counter im lokalen Array erstellen
        Local_iAnalyse[i] = 0;                                        // Nullen für spätere Addition
        //~ // printf("Local_iAnalyse[i]=%i \n", Local_iAnalyse[i]);
    }


//
//     Analyse


//     Analyse der 'Local_String' Daten
    for (i = 0; i < iGlobalSize-1; i+=2) {                            // Lokales Array 'Local_String' durchgehen
        for (j = 0; j < iNumAnalyse*2-1; j+=2) {                    // Und bei Erfolg 'Local_iAnalyse[i]++'
            if (Local_String[i] == Local_sAnalyse[j]) {
                if (Local_String[i+1] == Local_sAnalyse[j+1]) {
                    Local_iAnalyse[j/2]+=1;
                    //~ // printf("j/2=%i, Local_iAnalyse[j/2]=%i \n", j/2, Local_iAnalyse[j/2]);
                }
            }
        }
    }
    // Debug 'Local_iAnalyse'
    for (i = 0; i < iNumAnalyse; i++) {
        //~ // if (Local_iAnalyse[i] > 0) {printf("i=%i, Local_iAnalyse[i]=%i \n", i, Local_iAnalyse[i]);}
    }


//     Analyse-Ergebniss im globalen Speicher addieren
    barrier(CLK_GLOBAL_MEM_FENCE);                                     // Warten, bis alle Workitems fertig sind!
    for (i = 0; i < iNumAnalyse; i++) {                                // Ergebniss addieren
        if (Local_iAnalyse[i] > 0) {atomic_add(&iAnalyse[i], Local_iAnalyse[i]);}
        //~ // if (Local_iAnalyse[i] > 0) {iAnalyse[i] += Local_iAnalyse[i];}
        //~ //printf("iAnalyse[i]=%i \n", iAnalyse[i]);
    }
    // Übergebe ich die Daten falsch an das globale Array zurück ?!? ...


//     Fertiggestelltes Analyse-Ergebniss in 'Local_iNewAnalyse' kopieren
    barrier(CLK_GLOBAL_MEM_FENCE);                                     // Warten, bis alle Workitems fertig sind!
    for (i = 0; i < iNumAnalyse; i++) {                                // Analyse-Counter im lokalen Array erstellen
        Local_iNewAnalyse[i][0] = iAnalyse[i];                        // Ergebniss in lokal schreiben
        Local_iNewAnalyse[i][1] = i;                                // Index für Sortierung
        //~ // printf("Local_iNewAnalyse[i][0]=%i, Local_iNewAnalyse[i][1]=%i \n", Local_iNewAnalyse[i][0], Local_iNewAnalyse[i][1]);
    }
    // jedenfalls können die Ergebnisse nicht stimmen...
    // Sind das wieder Syncronisations-Probleme, oder habe ich was falsch gemacht??


//     'Local_iNewAnalyse' sortieren


//
//     Replacing


//     'Local_iNewAnalyse' tauschen mit (uInt)Chars welche nicht in 'Local_NotReplace' vorhanden sind


//
//     Rückgabe


}


/*
    SimpleCompression:
    Dieser Kernel soll eine minimal Kompression in OpenCL darstellen.
    Ich wandel alle Char´s in uInt um, ich denke das so die Grafikkarten evtl. besser Arbeiten können (Eine Vermutung)
    Momentan hänge ich noch an der Analyse, der Array 'iAnalyse' liefert noch unerklärliche Daten...
    Spätere Versionen sollen noch weitere Methoden zur Analyse enthalten
    Auch soll noch der DeCompress-Key, welcher zur Entpackung nötig ist, wie im AU3 Orginal noch verarbeitet werden
*/

Alles anzeigen

Zu der Vorgehensweise, dem Kernel wird ein Datencluster (25KB) übergeben, dieser wird intern in '$iGlobalSize' große Teile geteilt. Ich Teste den Kernel mit '_CL_RunKernel(1, 0)' um nur mit einem Thread zu Arbeiten. Aber wie weise ich jetzt die richtige globale Threadanzahl zu? muß ich '_CL_RunKernel(int(StringLen($sInput) / $iGlobalSize)+1, 0)' verwenden?? Momentan zickt der Kernel noch rum deswegen lässt sich das schlecht testen...
Erstmal werde ich versuchen diesen Kernel in dieser Weise weiter auszubauen, sobald dieser richtig lauffähig ist möchte ich noch 'get_local_id(0)' integrieren - so wie ich das verstanden habe hat jeder globaler Thread ja noch seine lokalen Einheiten. (Ich hoffe das ist jetzt nicht zu Fantastisch)
Bei 'atomic_add(&iAnalyse[$i], Local_iAnalyse[$i])' übergebe ich an die Funktion 'atomic_add' den Pointer auf dem Index des 'iAnalyse' Array´s. Ich möchte die Analyse- und die Replacefunktion in 'Void' Funktionen packen und dort eben das gesammte Array übergeben - mache ich das denn so '_OclAnalyse(&iAnalyse)' auf 'void _OclAnalyse(uint *iArray)' ?? <- Dabei wird dieses Array doch 'ByRef' übergeben, oder?

[Multiprocessing]
In dem Script 'SimpleCompression_MP.au3' muß man ab Zeile 27 seine entsprechenden OpenCL-Devies eintragen - war die schnellste Lösung... Die Prozesse müssen kompiliert werden, werden als *.exe aufgerufen. Der 'SimpleCompression_AU3' Prozess läuft schonmal wie erwartet, und der 'SimpleCompression_OCL' verwendet den gleichen Kernel wie in 'SimpleCompression.au3' - ist somit noch nicht lauffähig. Aber die Basis des Multiprocessing sollte Lauffähig sein, bei den OCL-Part konnte nicht richtig testen aber es gibt bei mir 3 Rückgaben welcher meinen Devices entspricht.

Grüsse!

[EDIT]
Solltet ihr seltsamme Erscheiningungen in der OpenCL installation haben, wie z.B. mehr als einen Prozessor - obwohl ihr nur einen installiert habt; Dann habt ihr wahrscheinlich bei Intel Systemen erstmal die Runtime und anschliessend das SDK installiert. Ihr müsst einfach nur die Runtime DeInstallieren dann klärt sich das Problem.

[EDIT2]
Ich habe vergessen in 'SimpleCompression_MP.au3' die Zeile 101 zu De-Kommentieren, wenn diese Aktiviert wird startet auch das AutoIt Multiprocessing...

**Andy**

Zitat von Techmix

Ich Teste den Kernel mit '_CL_RunKernel(1, 0)' um nur mit einem Thread zu Arbeiten.

Nicht mit einem Thread, sondern mit einem Workitem! Das ist ein Unterschied!

Zitat von Techmix

Aber wie weise ich jetzt die richtige globale Threadanzahl zu? muß ich '_CL_RunKernel(int(StringLen($sInput) / $iGlobalSize)+1, 0)' verwenden??

Schau mal in die References zum Thema Speicherverwaltung. Ggf. HIER und HIER die Grafiken.

Generell ist es aber unnötig, nur "ein" Workitem rechnen zu lassen, der Kernel läuft doch sowieso für alle anderen Workitems parallel. Ich verstehe jedenfalls deinen Ansatz nicht. Das spricht gegen sämtliche Ideen und Grundsätze von OCL!
"Global" kannst du den Speicherbereich so groß machen wie du willst, problematisch wird es nur mit der Aufteilung auf die Workgroups und die Workitems. Die Aufteilung auf Workgroups und Workitems übernimmt doch sowieso OpenCL (abhängig vonn der Hardware) für dich!
WENN du lokalen Speicher (innerhalb der Workitems __PRIVATE, innerhalb der Workgroups __LOCAL, Workgroups synchronisieren NUR über __GLOBAL (langsam) oder __CONSTANT) benutzen willst/musst, pass auf, denn je nach Hardware sind dort nur 16 bzw. 32K Speicher vorhanden! (Abfrage bei den Devices)

Zitat von Techmix

Dabei wird dieses Array doch 'ByRef' übergeben, oder?

Man kann keine Arrays an den OCL-Kernel übergeben. Schreib das Array in die Struct, übergib den Pointer. Ob du dann diesen Speicher überschreibst (das meinst du hoffentlich mit ByRef) oder die Ergebnisse in einem weiteren Speicherbereich zurückgibst, ist unerheblich!
Du machst dir dein Leben auch dermaßen schwer^^
Ich würde ganz pragmatisch Stückchen für Stückchen vorgehen. Daten an den Kernel übergeben und mir dann die Zwischenergebnisse zurückgeben lassen.

Ich habe immer noch das Gefühl, dass du dir viel zu viel Gedanken über die Aufteilung innerhalb der Workitems machst!
Wie schon in den ersten Posts hier gesagt, schreib EINFACHE Funktionen bzw. Kernel! Ich verstehe ehrlich gesagt dieses gesamte rumgefriemel mit den Kernelparametern und AutoItvariablen nicht! Optimiert wird grundsätzlich erst dann, wenn das Programm läuft! Und dann ist es WESENTLICH effektiver, erst einmal die Optimierungsmöglichkeiten innerhalb von OCL auszuschöpfen, als immer die SLI-Variante im Hinterkopf zu haben.

Zitat von Techmix

Aber wie weise ich jetzt die richtige globale Threadanzahl zu?

Darauf solltest du dich konzentrieren...
Ich habe mal Infos zu einer Radeon 7970 gefunden:

Zitat

The AMD Radeon HD 7970 is one of the AMD Southern Islands series that also
adopts the GCN architecture. The HD 7970 contains 32 compute units (CUs).
All CUs share a 768 KB L2 cache. Each CU has four vector units, a scalar unit,
a 16 KB read/write L1 cache, and a 64 KB local data share (32 KB is available
to a single wavefront). Each vector unit consists of 16 processing elements (PEs);
thus, the HD 7970 is equipped with 2048 (= 32 * 4 * 16) PEs

Du hast also 32 Compute Units a 4 Vector Units a 16 Processing Elements (das sind die "Workitems").
Man sollte also schauen, dass möglichst viele dieser Workitems gleichzeitig "Futter" haben. Im Endeffekt ist es kontraproduktiv, wenn du bei der 7970 bspw. 2,2MB Daten in je 1000Bytes großen Blöcken übergeben würdest. Dann würden 2048 Workitems parallel abgearbeitet, und die restlichen 152 "danach". Das heisst, 1896 Processing units laufen LEER und warten darauf, dass die 152 PE´s fertig werden!!!!
Daher berechnet man entweder anhand seiner Hardware die Blockgröße, oder lässt (wie schon mehrmals gezeigt) einfach verschiedene Blockgrößen durchlaufen und merkt sich diejenige mit der schnellsten Berechnungszeit! Das allerdings ist das LETZTE, was bei der OCL-Programmierung in Angriff genommen werden sollte...

EDIT Typo!!!!!!! Muss
$CL_iAnalyse = _CL_CreateBuffer($p_iAnalyse)

Techmix

Hallo

Zitat

...Ich verstehe jedenfalls deinen Ansatz nicht...

Das mit dem einen Workitem mache ich nur, damit ich die printf() Ausgaben besser analysieren kann! Ist quasi NUR zu Debbugzwecken gesetzt, wird natürlich noch angepasst...

Zitat

...das meinst du hoffentlich mit ByRef...

Ich meine das innerhalb des Kernels (der *.cl Datei), wenn ich aus dem Kernel eine eigene Funktion aufrufe und ich eben ein Array als Rückgabe brauche. In anderen Sprachen muss ich "Arrays" ByRef an die Funktion übergeben - weil es nicht möglich ist ein Array aus der Funktion zurückgeben zu lassen

Zitat

Du machst dir dein Leben auch dermaßen schwer

Grins... Ich Denke einfach zu kompliziert, kann mich aber nicht entsprechend äussern/ausdrücken... DAS ist das Problem!

Zitat

...Kernelparameter und AutoItvariablen...

Ganz einfach, in OpenCL kann ich keine variablen Arrays verwenden - diese müssen eine vordefinierte Größe erhalten. Aber je nach Datengröße müsste ich immer den Kernel händisch anpassen, jetzt nicht mehr! Das übernimmt AU3 für mich... UND ich befinde mich doch in der Lernphase dieses Rumgeieiere erleichtert mir einfach den Umgang...

OK, der Kernel befindet sich noch in der Analyse von dem Input. Ich erstelle in AU3 eine 'uInt' Struct, Nulle diese und übergebe sie mit '__global uint* iAnalyse' an den Kernel. Am Kernel angekommen lasse ich mir dieses mit printf() ausgeben, und es ist nicht mehr 0. Die komischten Werte werden mir ausgespuckt, mal eine Null, sogar negative Zahlen werden ausgespuckt. Dabei erstelle ich die Struct wie alle anderen auch, und übergebe diese auch ganz normal. Alle anderen Structs kann ich mir ebenso mit printf() ausgeben, und da sind die Daten wie erwartet... Liegt das jetzt an einer unachtsammkeit innerhalb meines AU3 Scriptes, oder an meinem Kernel selber

Ausserdem finde ich meinen Kernel garnicht soooo kompliziert... Der Vorgang liegt ersichtlich im AU3 Source, ist eine einfache Analyse der Eingangsdaten, mit anschliessender Sortierung der Ergebnisse und Austausch mit freien Chars. Anschliessend alles in die Output-Struct schreiben und in AU3 den Buffer auslesen - Und fertig! <- Wo ist das kompliziert

Ich habe das Multiprocessing absichtlich eingebaut, das ist auch nicht wirklich aktuell hab das schon vor 2 Jahren verwendet. In manchen Bereichen kann ich damit AutoIt einen richtigen Turbo verpassen. Und dieses lässt sich anhand von dieser Billig-Kompression sehr schön darstellen, im normalen Singleprozess-Betrieb braucht das standart Script 63 Sekunden bei mir um die 'AutoIt.exe' zu komprimieren, im Multiprozess-Betrieb nur noch 18 Sekunden - das ist drei mal schneller als normal. Und wieso nicht mit OpenCL ebenso vorgehen - es bietet sich mit den unterschiedlichen Platformen und Devices doch dafür an . Mir geht es auch nicht um SLI oder 'prestige' (hier guck mal, Boahhh...) für mich ist das eine normale Vorgehensweise! Maximierung der Rechenleistung durch Auslagerung in Prozesse oder Threads
Daher bin ich auch schon so angefressen, das ich mich mit den Fortschritt in OpenCL "etwas" schwer tue! Immer wieder Stolpersteine, die OCL SDK´s lassen sich nicht in die 'Express' Versionen vom VisualStudio integrieren, dann das dumm gelaufene Update, ... Immerhin konnte ich SciTE überreden mir endlich etwas Farbe im Kernel anzeigen zu lassen

Grüsse!

[Edit]
Hab jetzt erst dein Edit gelesen... Das ist ja mal peinlich Aber Danke!

**Andy**

Zitat von Techmix

Die komischten Werte werden mir ausgespuckt, mal eine Null, sogar negative Zahlen werden ausgespuckt.

Das hat mich bei meiner ersten Initialisierung eines Arrays im lokalen Speicher auch gefrustet. Das Problem war, dass OCL wohl die Arrays nicht "nullt", also mit Nullen füllt! Da wird einfach ein Speicherbereich reserviert, mit allem, was darin steht^^. Also sicherheitshalber das Array "nullen".

Zitat von Techmix

...Kernelparameter und AutoItvariablen...

Ganz einfach, in OpenCL kann ich keine variablen Arrays verwenden - diese müssen eine vordefinierte Größe erhalten. Aber je nach Datengröße müsste ich immer den Kernel händisch anpassen, jetzt nicht mehr! Das übernimmt AU3 für mich... UND ich befinde mich doch in der Lernphase dieses Rumgeieiere erleichtert mir einfach den Umgang...

Die Idee finde ich klasse! Sollten wir weiter ausbauen! Da gibts bestimmt noch weitere "Tricks" die man in die Kernel integrieren könnte.
Wusstest du übrigens, dass man außerhalb vom Kernel Funktionen schreiben kann?

Zitat von Techmix

Daher bin ich auch schon so angefressen, das ich mich mit den Fortschritt in OpenCL "etwas" schwer tue! Immer wieder Stolpersteine, die OCL SDK´s lassen sich nicht in die 'Express' Versionen vom VisualStudio integrieren, dann das dumm gelaufene Update, ... Immerhin konnte ich SciTE überreden mir endlich etwas Farbe im Kernel anzeigen zu lassen

Immer schön geschmeidig bleiben...
Manch anderer wäre froh, wenn er dieses KnowHow hätte!
Ich selbst bin nur rudimentär in C (und somit den OCL-Kerneln) fit, leider.
Man müsste sich viel mehr damit beschäftigen.

Schaumamal wo das hinführt, ich habe ein gutes Gefühl^^. Gerade mit dem Hintergrund, dass eben NICHT dieser Mega-Supertanker Visual Studio für die Arbeit mit OpenCL nötig ist, sondern eine "billige" Scriptsprache

Techmix

Hy,
Wie bekomme ich es mit einigen Zeilen so hin, das ich an einem CL-Kernel einen Synaxcheck ausführen kann? Mir geht es um diese SciTE Ausgabe:

Code

LOG:
Compilation started
1:33:29: error: parameter may not be qualified with an address space
1:33:73: error: parameter may not be qualified with an address space
...

Ich habe doch noch 'OpenCL4SciTE' erstellt, habe da ein kleines Script für ein Helpfile integriert - Jetzt möchte ich noch den Kernel direkt aus SciTE kompilieren um damit einen Synaxcheck zu integrieren. Es nervt mich das ich andauernd in´s AU3 Script wechseln muss um im Kernel einen Synatax-Check durchführen zu können!
In der Funktion '_cl_GetDevice()' wird das Logfile erstellt, auf welches ich so scharf bin Nur leider wird zum Erstellen des Log´s großes PiPaPo betrieben, und ich bin mir nicht sicher wie oder welche Teile davon ich verwenden soll...

[OpenCL4SciTE]
Auch stimmt was nicht in der 'opencl.properties' nicht, wieso wird das API-File nicht verwendet?!
Mir ist auch nicht entgangen das nicht alle Keywords verwendet werden, habe schon einige manuell zugefügt. Das ist einfach passiert weil ich mir ein Highlighting von einem anderen Editor aus dem Netz geladen habe und diese Daten eben dür SciTE´s Format angepasst.
Auch hätte ich die Farben in SciTE näher an AU3, also keine roten Funktionsbezeichnungen oder lila Textinhalt - sondern Hellblaue Funktionen und grauer Text, da stimmen sicherlich noch andere Farben nicht aber das ist störend. Ich habe schon mit den 'style' Einstellungen gespielt, muß ich diese wirklich händisch sortieren oder reicht eine richtige 'keywordsNR' Zuweisung - obwohl ich auch hier experimentiret habe. Auch macht es einen riesen Unterschied ob ich bei 'lexer.$(ocl)=cpp' oder einfach was anderes angebe. So viele Möglichkeiten, da steige ich nicht ganz durch...
Ich hänge noch die aktuelle Datei dazu an, enthällt kleinere Änderungen...

Grüsse!

**Andy**

Zitat von Techmix

In der Funktion '_cl_GetDevice()' wird das Logfile erstellt, auf welches ich so scharf bin Nur leider wird zum Erstellen des Log´s großes PiPaPo betrieben, und ich bin mir nicht sicher wie oder welche Teile davon ich verwenden soll...

Ich nehme das mal in Angriff...
Übrigens bekomme ich von der AMD/Radeon-Implementierung bspw folgende Fehler/Warnungen

Zitat

LOG:
"C:\Users\Mama\AppData\Local\Temp\OCL966B.tmp.cl", line 6: warning: variable
"p" was declared but never referenced
uint i,r,p=divisor;
^

Bekommst du nur die von dir geposteten "Fehlermeldungen"?

Techmix

Ja, diese ist mir fremd. Ausserdem welche Variable p?? Habe diese jedenfalls nicht deklariert...

[Edit]
Das sieht mir danach aus, als würde der Compiler dazwischen funken...

**Andy**

nein, die Variable p hatte ICH eingebaut, um die Warnung zu provozieren
Und ja, der Compiler wirft diese Meldungen!
Sehr aussagekräftig jedenfalls!
Was hälst du davon, wenn als Rückgabe einer Funktion _CL_CheckKernelSyntax($kernel) das Logfile zurückgegeben wird. Im optimalen Fall (kein Fehler/Warnung) wird ein Leerstring zurückgegeben, ansonsten das Log. Ggf. sollte man drüber nachdenken, die Warnungen auch auszublenden und nur bei "echten" Fehlern das Log zurückzugeben.

Techmix

Das wäre super! Würde den Umgang mit den Kerneln deutlich erleichtern. Ich halte das eh für eine gute Idee das in SciTE zu integrieren weil die meißten AU3 User sowieso den Editor verwenden. Und wie du sagtest, so kann man besser VS umgehen und effizienter Arbeiten.

Grüsse!

[Edit]
Die warnungen haben mich anfangs auch sehr irritiert, habe versucht diese "Wegzuarbeiten"... Können ruhig ausbleiben!

**Andy**

hehe, wie wärs, beim Editieren des Kernels auf die F5-Taste zu hämmern und in der Console steht das Log oder "Kernel ohne Syntax-Fehler compiliert!"

Techmix

Genau so hatte ich mir das vorgestellt, habe auch schon die "opencl.properties"entsprechend vorbereitet da hab ich schon ein .chm Helpfile auf F1 integriert.

Techmix

Hallo ich hab wieder einen Hängerchen

Es geht voran, aber ich habe gerade ein Problemchen mit einer Poolzuweisung. Obwohl ich im Kernel 'if (Local_NotReplace[j] != i)' verwende wird an dieser Position das Ungleich '!=' völlig übergangen und das resultierende Array wir nicht korreckt aufgebaut. UND ich bin endlich bei der letzten großen Funktion angekommen, den Replacer (ab Zeile 184). Dort benötige ich noch das Pool-Array welches nicht richtig erstellt wird, habe aber noch Probleme die Durchlauflogik richtig umzusetzen - Es soll der Pool überprüft werden, und ein vorhandensein des aktuellen Pool-Chars im 'String' (In AU3 würde ich StringInString() verwenden). Wenn erst das Pool-Array richtig erstellt wird, denke ich das der Rest nicht mehr allzuviel Schwierigkeiten machen wird

Du erwähntest das man im *.cl Kernel auch Funktionen ausserhalb des __Kernel erstelen und ausführen kann. Ich habe eine Analyse- und eine Replacefunktion bereits in Void Funktionen eingebaut. Die identischen Durchläufe befinden sich jetzt auch noch zusätzlich im __Kernel, weil mir nur Fehlermeldungen ausgespuckt werden egal wie ich in die Funktionen springe, oder wie ich die Parameter übernehme (Zeile 148 bis 162). Deswegen habe ich auch die beiden Funktionen Auskommentiert...

Grüsse!

CL-Kernel

Code

/*
    SimpleCompression:


    Master:
    Im Master-Script werden die Eingangsdaten in Daten-Cluster aufgeteilt (Momentan 25KB, muß noch optimiert werden)
    Die Daten-Cluster werden nacheinander an den OpenCL-Kernel übergegeben
    In AU3 werden Variablen ausgetauscht um Arraygrößen im Kernel flexibel zuweisen zu können


    Kernel-StartUP:
    Der Kernel teilt seine Eingangsdaten in '$iGlobalSize' große Teile in 'Local_String' auf - Umwandlung dabei nach uInt (wie bei Asc())
    Für alle wichtigen Arrays werden 'Lokale' erstellt und reinkopiert


    Analyse:
    Char Wiederholungen in 'Local_iNewAnalyse' zählen
    Analyse von jeweils 2 'Local_String's, wenn nicht in 'NotReplace' vorhanden 'Local_iAnalyse[i]+=1'
    Analyse-Ergebniss in globalen Speicher addieren
    Fertiggestelltes Analyse-Ergebniss in 'Local_iNewAnalyse' zurück kopieren


    Replacing:
    'Local_iNewAnalyse' absteigend sortieren
    ? 'Local_iNewAnalyse' tauschen mit (uInt)Chars welche nicht in 'Local_NotReplace' vorhanden sind


    Rückgabe:
    ???


*/


//~ void _Sorter(private uint* _iNewAnalyse, uint iNumAnalyse);
//~ void _Replacer(private uint _Input, private uint* _Output, private uint _Array, private uint _NotReplace, private uint* _Replaced);


#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable    // Atomics einschalten


__kernel void SimpleCompression(
                __constant char* _In,                                 // Input
                __global char* _Out,                                // Output - Rückgabe
                __constant uint* NotReplace,                        // NotReplace
                __constant char* Analyse,                             // Analyse-Chars
                __global uint* iAnalyse)                            // Analyse-Counter
{
// Deklarationen
    uint iTrenner        = NotReplace[0];                            // Trenner, für die Schlüsselerzeugung
    uint iNumElements    = $iNumElements;                            // Input Länge ($iNumElements wird in AU3 Replaced)
    uint iNumNotReplace    = $iNumNotReplace;                            // Anzahl von Not-Replacern ($iNumNotReplace wird in AU3 Replaced)
    uint iGlobalSize    = $iGlobalSize;                                // Eingangs-Datengröße für Globale Daten ($iGlobalSize wird in AU3 Replaced)
    uint iLocalSize        = $iLocalSize;                                // Datengröße je Thread für Lokale Daten ($iLocalSize wird in AU3 Replaced)
    uint iNumAnalyse    = $iNumAnalyse;                                // Anzahl von Analyse-Einträgen ($iNumAnalyse wird in AU3 Replaced)
    uint iThreadID        = get_global_id(0)*iGlobalSize;                // Global-ThreadID alle iDataSize-Char
    uint iLocalID        = get_local_id(0)*2;                        // Local-ThreadID alle 2-Char
    uint iOutput        = 0;                                        // Output Counter
    uint i, j;


// Lokale Array´s
    private uint Local_String[$iGlobalSize];                        // Lokales Array 'Local_String'        ($iGlobalSize wird in AU3 Replaced)
    private char Local_sOutput[$iGlobalSize];                        // Lokales Array 'Local_sOutput'    ($iGlobalSize wird in AU3 Replaced)
    private uint Local_NotReplace[$iNumNotReplace];                    // Lokales Array 'Local_NotReplace'    ($iNumNotReplace wird in AU3 Replaced)
    private uint Local_sAnalyse[$iNumAnalyse*2];                    // Lokales Array 'Local_sAnalyse'    ($iNumAnalyse wird in AU3 Replaced)
    private uint Local_iAnalyse[$iNumAnalyse];                        // Lokales Array 'Local_iAnalyse'
    private uint Local_iNewAnalyse[$iNumAnalyse][2];                // Lokales Array 'Local_iNewAnalyse'
    private uint Local_iReplaced[$iNumAnalyse][3];                    // Lokales Array 'Local_iReplaced'    Schlüsselerzeugung
    private uint Local_iPool[$iNumAnalyse];                            // Lokales Array 'Local_iPool'        Freie Replacer


// Debugging
    //~ printf("get_global_id(0)=%i \n", get_global_id(0));            // 0
    //~ printf("get_global_size(0)=%i \n", get_global_size(0));        // 1
    //~ printf("get_local_id(0)=%i \n", get_local_id(0));            // 0
    //~ printf("get_local_size(0)=%i \n", get_local_size(0));        // 1
    //~ printf("iGlobalSize=%i \n", iGlobalSize);
    //~ printf("iNumElements=%i \n", iNumElements);
    //~ printf("iNumNotReplace=%i \n", iNumNotReplace);
    //~ printf("iNumAnalyse=%i \n", iNumAnalyse);


//
//     Kernel-StartUP


//     Input-Daten verarbeiten...
    for (i = 0; i < iGlobalSize; i++) {                                // 'iGlobalSize' Chars von 'String' in lokales Array kopieren
        if (iThreadID+i < iNumElements) {                            // Aufpassen das 'iThreadID+i' < 'iNumElements' bleibt
            Local_String[i] = convert_int(_In[iThreadID+i]);        // Umwandlung von Char zu Int
            //~ printf("iThreadID+i=%i, _In[iThreadID+i][i]=%c, Local_String[i]=%i \n", iThreadID+i, _In[iThreadID+i], Local_String[i]);
        }
    }
    for (i = 0; i < iGlobalSize; i++) {                                // 'Local_sOutput' Leeren
        Local_sOutput[i] = "";
    }
    for (i = 0; i < iNumNotReplace; i++) {                            // NotReplace in lokales Array kopieren
        Local_NotReplace[i] = NotReplace[i];
        //~ printf("Local_NotReplace[i]=%i \n", Local_NotReplace[i]);
    }
    for (i = 0; i < iNumAnalyse*2; i++) {                            // Analyse in lokales Array kopieren
        Local_sAnalyse[i] = convert_int(Analyse[i]);
        //~ printf("Local_sAnalyse[i]=%i, Local_sAnalyse[i]=%c \n", Local_sAnalyse[i], Local_sAnalyse[i]);
    }
    for (i = 0; i < iNumAnalyse; i++) {                                // Analyse-Counter im lokalen Array erstellen
        Local_iAnalyse[i] = 0;                                        // Nullen für spätere Addition
    }
    for (i = 0; i < iNumAnalyse; i++) {                                // Zweiten Analyse-Counter im lokalen Array erstellen
        Local_iNewAnalyse[i][0] = 0;                                // Nullen
        Local_iNewAnalyse[i][1] = 0;                                // Nullen
    }
    for (i = 0; i < iGlobalSize; i++) {                                // Output im lokalen Array erstellen
        Local_sOutput[i] = "";                                        // Leeren für spätere Aktion
    }
    for (i = 0; i < iNumAnalyse; i++) {                                // Schlüssel-Array Erstellung
        Local_iReplaced[i][0] = -1;                                    // Replacer
        Local_iReplaced[i][1] = -1;                                    // Analyse Char 1
        Local_iReplaced[i][2] = -1;                                    // Analyse Char 2
    }
    for (i = 32; i < iNumAnalyse; i++) {                            // Pool-Array Erstellung
        Local_iPool[i] = -1;                                        // Leeren
        for (j = 0; j < iNumNotReplace; j++) {                        // Prüfen ob 'i' in 'Local_NotReplace' enthalten ist
            if (Local_NotReplace[j] != i) {
        //~ printf("Local_NotReplace[j]=%i \n", Local_NotReplace[j]);
                Local_iPool[i] = i;                                    // Dann 'Local_iPool[i]' = 'i'
            }
        }
        printf("Local_iPool[i]=%i \n", Local_iPool[i]);
    }
//~ private uint Local_iPool[$iNumAnalyse];


//
//     Analyse


//     Analyse der 'Local_String' Daten
    for (i = 0; i < iGlobalSize-1; i+=2) {                            // Lokales Array 'Local_String' durchgehen
        for (j = 0; j < iNumAnalyse*2-1; j+=2) {                    // Und bei Erfolg 'Local_iAnalyse[i]++'
            if (Local_String[i] == Local_sAnalyse[j]) {
                if (Local_String[i+1] == Local_sAnalyse[j+1]) {
                    Local_iAnalyse[j/2]++;
                    //~ printf("j/2=%i, Local_iAnalyse[j/2]=%i \n", j/2, Local_iAnalyse[j/2]);
                }
            }
        }
    }
    // Debug 'Local_iAnalyse'
    //~ for (i = 0; i < iNumAnalyse; i++) {
        //~ if (Local_iAnalyse[i] > 0) {printf("i=%i, Local_iAnalyse[i]=%i \n", i, Local_iAnalyse[i]);}
    //~ }


//     Analyse-Ergebniss im globalen Speicher addieren
    barrier(CLK_GLOBAL_MEM_FENCE);                                     // Warten, bis alle Workitems fertig sind!
    for (i = 0; i < iNumAnalyse; i++) {                                // Ergebniss addieren
        if (Local_iAnalyse[i] > 0) {atomic_add(&iAnalyse[i], Local_iAnalyse[i]);}
        //~ if (Local_iAnalyse[i] > 0) {iAnalyse[i] += Local_iAnalyse[i];}
        //~ printf("iAnalyse[i]=%i \n", iAnalyse[i]);
    }


//     Fertiggestelltes Analyse-Ergebniss in 'Local_iNewAnalyse' kopieren
    barrier(CLK_GLOBAL_MEM_FENCE);                                     // Warten, bis alle Workitems fertig sind!
    for (i = 0; i < iNumAnalyse; i++) {                                // Analyse-Counter im lokalen Array erstellen
        Local_iNewAnalyse[i][0] = iAnalyse[i];                        // Anzahl in lokal schreiben
        Local_iNewAnalyse[i][1] = i;                                // uIntChar für Sortierung
        //~ printf("Local_iNewAnalyse[i][0]=%i, Local_iNewAnalyse[i][1]=%i \n", Local_iNewAnalyse[i][0], Local_iNewAnalyse[i][1]);
    }


//
//     Replacing


//     'Local_iNewAnalyse' sortieren
    //~ _Sorter(Local_iNewAnalyse, iNumAnalyse);
    uint iWert1, iWert2, iElemente = iNumAnalyse;
    barrier(CLK_LOCAL_MEM_FENCE);                                    // Lokaler Barrier
    while(iElemente--){                                                // BubbleSort2D, angepasst aus C
        for(i = 1; i <= iElemente; i++){
            if(Local_iNewAnalyse[i-1][0] < Local_iNewAnalyse[i][0]){
                iWert1 = Local_iNewAnalyse[i][0];                    // Anzahl
                iWert2 = Local_iNewAnalyse[i][1];                    // uIntChar
                Local_iNewAnalyse[i][0] = Local_iNewAnalyse[i-1][0];
                Local_iNewAnalyse[i][1] = Local_iNewAnalyse[i-1][1];
                Local_iNewAnalyse[i-1][0] = iWert1;
                Local_iNewAnalyse[i-1][1] = iWert2;
            }
        }
    }
    // Debug sortiertes 'Local_iNewAnalyse'
    //~ for (i = 0; i < iNumAnalyse; i++) {                                // Analyse-Counter im lokalen Array erstellen
        //~ printf("Sorted: Local_iNewAnalyse[i][0]=%i, Local_iNewAnalyse[i][1]=%i \n", Local_iNewAnalyse[i][0], Local_iNewAnalyse[i][1]);
    //~ }


//     'Local_iNewAnalyse' tauschen mit (uInt)Chars welche nicht in 'Local_NotReplace' vorhanden sind
    //~ _Replacer(Local_String, Local_sOutput, Local_iNewAnalyse, Local_NotReplace, Local_iReplaced);
    uint uIntChar1, uIntChar2, iCount = 0;
    //~ for (i = 0; i < iGlobalSize; i++) {
        //~ printf("Local_sOutput[i]=%s \n", Local_sOutput[i]);
    //~ }
    barrier(CLK_LOCAL_MEM_FENCE);                                    // Lokaler Barrier
    for (i = 0; i < iGlobalSize-1; i+=2) {
        uIntChar1 = Local_String[i];                                // uIntChar´s erstellen
        uIntChar2 = Local_String[i+1];
        for (j = 0; j < iNumAnalyse; j+=2) {                        // Lokales Array 'Local_String' durchgehen
            if (Local_iNewAnalyse[j][0] > 1) {
                if (Local_String[i] == Local_sAnalyse[Local_iNewAnalyse[j][1]]) {
                    if (Local_String[i+1] == Local_sAnalyse[Local_iNewAnalyse[j][1]+1]) {
                    //~ if (Local_String[i+1] == Local_sAnalyse[Local_iNewAnalyse[j][1]+1]) {
                        //~ private uint Local_NotReplace[$iNumNotReplace];
                        //~ private uint Local_iReplaced[$iNumAnalyse][3];
                        uIntChar1 = Local_iNewAnalyse[j][1];
                        uIntChar2 = -1;
                    }
                }
            }
        }
        if (iCount < iNumElements) {                                // Aufpassen das 'iCount' < 'iNumElements' bleibt
            Local_sOutput[iCount] = convert_char(uIntChar1);        // uIntChar1 schreiben
            iCount++;
            if (uIntChar2 != -1){                                    // uIntChar2 schreiben, wenn nicht Leer
                if (iCount < iNumElements) {                        // Aufpassen das 'iCount' < 'iNumElements' bleibt
                    Local_sOutput[iCount] = convert_char(uIntChar2);
                    iCount++;
                }
            }
        }
    }
    for (i = 0; i < iGlobalSize; i++) {
        printf("Local_sOutput[i]=%c \n", Local_sOutput[i]);
    }


//
//     Rückgabe erstellen


//    '_Out' mit 'Local_sOutput' befüllen
    j = 0;
    barrier(CLK_GLOBAL_MEM_FENCE);                                     // Warten, bis alle Workitems fertig sind!
    for (i = 0; i < iGlobalSize; i++) {                                // Lokales Array 'Local_sOutput' durchgehen
        if (Local_sOutput[i] != ""){
            _Out[j] = Local_sOutput[i];
            j++;
        }
    }


//    Schlüsselerzeugung
    barrier(CLK_GLOBAL_MEM_FENCE);                                     // Warten, bis alle Workitems fertig sind!
    for (i = 0; i < iNumAnalyse; i++) {                                // Schlüssel anhängen
        if (Local_iReplaced[i][0] > -1){
            _Out[j] = convert_char(iTrenner);                        // Trenner
            _Out[j+1] = convert_char(Local_iReplaced[i][0]);        // Replacer
            _Out[j+2] = convert_char(iTrenner);                        // Trenner
            _Out[j+3] = convert_char(Local_iReplaced[i][1]);        // Char 1
            _Out[j+4] = convert_char(Local_iReplaced[i][2]);        // Char 2
            j+=5;
        }
    }


}


// Schlüsselerstellung in AU3
//~ $sData = $sString & $a_Trenner[0][1]


//~ ; $sKey Erstellen
//~ For $i = 0 To UBound($aReplaced, 1) - 1
    //~ If $aReplaced[$i][0] <> "" And $aReplaced[$i][1] <> "" Then $sKey &= $aReplaced[$i][0] & $a_Trenner[0][1] & $aReplaced[$i][1] & $a_Trenner[0][1]
//~ Next
//~ $sData &= StringTrimRight($sKey, 1)


//~ void _Sorter(                                                        // Sortier Funktion
        //~ private uint* _iNewAnalyse,                                 // Analyse Array - Rückgabe
        //~ uint iElemente)                                                // Größe
//~ {
    //~ uint i, Wert1, Wert2;
    //~ barrier(CLK_LOCAL_MEM_FENCE);                                    // Lokaler Barrier
    //~ while(iElemente--){                                                // BubbleSort2D, angepasst aus C
        //~ for(i = 1; i <= iElemente; i++){
            //~ if(_iNewAnalyse[i-1][0] < _iNewAnalyse[i][0]){
                //~ Wert1 = _iNewAnalyse[i][0];                            // Anzahl
                //~ Wert2 = _iNewAnalyse[i][1];                            // uIntChar
                //~ _iNewAnalyse[i][0] = _iNewAnalyse[i-1][0];
                //~ _iNewAnalyse[i][1] = _iNewAnalyse[i-1][1];
                //~ _iNewAnalyse[i-1][0] = Wert1;                        // Neue Werte schreiben
                //~ _iNewAnalyse[i-1][1] = Wert2;
            //~ }
        //~ }
    //~ }
//~ }


//~ void _Replacer(                                                        // Austausch Funktion
                //~ private uint _Input,                                 // Eingangsdaten
                //~ private uint* _Output,                                 // Ausgang - Rückgabe
                //~ private uint _Array,                                 // _Sorter() Analyse-Array
                //~ private uint _NotReplace,                            // NotReplace
                //~ private uint* _Replaced)                            // Austausch-Array, benötigt zur Schlüsselerzeugung - Rückgabe
//~ {
    //~ uint uIntChar1, uIntChar2, iCount = 0;
    //~ barrier(CLK_LOCAL_MEM_FENCE);                                    // Lokaler Barrier
    //~ for (i = 0; i < iGlobalSize-1; j+=2) {
        //~ uIntChar1 = Local_String[i];                                // uIntChar´s erstellen
        //~ uIntChar2 = Local_String[i+1];
        //~ if (Local_iNewAnalyse[i][0] > 0) {
            //~ for (j = 0; j < iNumAnalyse; j+=2) {                        // Lokales Array 'Local_String' durchgehen
                //~ if (Local_String[i] == Local_sAnalyse[Local_iNewAnalyse[i][1]]) {
                    //~ if (Local_String[i+1] == Local_sAnalyse[Local_iNewAnalyse[i][1]+1]) {
                        //~ uIntChar1 = Local_iNewAnalyse[i][1];
                        //~ uIntChar2 = "";
                    //~ }
                //~ }
            //~ }
        //~ }
        //~ Local_sOutput[iCount] = uIntChar1;                            // uIntChar1 schreiben
        //~ iCount++;
        //~ if (uIntChar2 > ""){                                        // uIntChar2 schreiben, wenn nicht Leer
            //~ Local_sOutput[iCount] = uIntChar2;
            //~ iCount++;                                                // dann auch iCount++
        //~ }
    //~ }


//~ }

Alles anzeigen

**Andy**

Funktionen mit #define

Spoiler anzeigen

#define ORDERV(x,a,b) { bool swap = reverse ^ (getKey(x[a])<getKey(x[b])); \
data_t auxa = x[a]; data_t auxb = x[b]; \
x[a] = (swap)?auxb:auxa; x[b] = (swap)?auxa:auxb; }
#define B2V(x,a) { ORDERV(x,a,a+1) }
#define B4V(x,a) { for (int i4=0;i4<2;i4++) { ORDERV(x,a+i4,a+i4+2) } B2V(x,a) B2V(x,a+2) }
#define B8V(x,a) { for (int i8=0;i8<4;i8++) { ORDERV(x,a+i8,a+i8+4) } B4V(x,a) B4V(x,a+4) }
#define B16V(x,a) { for (int i16=0;i16<8;i16++) { ORDERV(x,a+i16,a+i16+8) } B8V(x,a) B8V(x,a+8) }

Zitat von Techmix

Obwohl ich im Kernel 'if (Local_NotReplace[j] != i)' verwende wird an dieser Position das Ungleich '!=' völlig übergangen und das resultierende Array wir nicht korreckt aufgebaut.

Puuuhhh, ich weiss ehrlich gesagt nicht, ob und wie != in OCL verwendet wird^^

Techmix

Hy,
Das Problem bei den "OpenCL Funktionen" ist, das ich versuchte ein 2D-Array zu übergeben. Hab das jetzt auf 2x1D-Array umgebaut, und ich kann OCL-Source in Funktionen auslagern.

Und ich habe in AU3 das "NotReplace" Array in ein "OnlyReplace" Array umgebaut, und kann dieses jetzt erfolgreich in OCL mit "==" statt "!=" abfragen

Grüsse!

**minx**

Brauche hier mal den Anhang

sedas19

würde mich auch sehr freuen wen der Anhang reupload wird, da der link nicht funktioniert

**Andy**

Wenn das mit dem Anhang nicht funktioniert, müsste ich eine neue ZIP bereitstellen.
Kann ich aber erst am Wochenende machen. Dann bitte nochmal melden

**Andy**

Hi zusammen,
aus aktuellem Anlass habe ich die OpenCL-Funktionen weiter angepasst und aufgefrischt.

Bitte die Beispielscripte testen, ggf. auch die Ergebnisse der Devices.au3 posten. Danke!

Es sollten die 64-Bit-Versionen der Scripte problemlos laufen, Geschwindigkeitsunterschiede zu 32Bit sind (bei mir zumindest) keine feststellbar,

Getestet wurde auf einem Intel-Notebook mit frisch installiertem und auf die neueste Version geupdateten Win10.
Im Prozessor ist ein Grafikchip integriert, dazu auf dem Board zusätzlich eine Nvidia-GPU. Die Tests liefen ohne Treiberupdates problemlos sowohl auf der CPU als auch auf den GPU´s.

Ich habe ein "Auswahlmenü" zur Verfügung gestellt, um das Gerät, auf welchem der Code ausgeführt werden soll, auswählen zu können.

Zwischenablage01.jpg
Dazu im Vergleich, meine "alte" Kiste