Problem mit OpenCL

[Nakai]

Heyho für genauere Problembeschreibung hier:
http://www.forum-3dcenter.org/vbulletin/showthread.php?t=537968

Code:
#define WORKGROUPSIZE 64
__kernel void MinMax(
__global float* input,
__global float * output,
float min,
float max,
int sizex,
int sizey,
int sizez
)
{
int idx = get_global_id(0);
int idy = get_global_id(1);
int idz = get_global_id(2);

int idx_local = get_local_id(0);
int idy_local = get_local_id(1);
int idz_local = get_local_id(2);

int local_x = get_local_size(0);
int local_y = get_local_size(1);
int local_z = get_local_size(2);

int group_x = get_num_groups(0);
int group_y = get_num_groups(1);
int group_z = get_num_groups(2);

int groupid_x = get_group_id(0);
int groupid_y = get_group_id(1);
int groupid_z = get_group_id(2);

int gid = idx + sizex*idy + sizex*sizey*idz;
int gid_local = idx_local + local_x*idy_local +local_x*local_y*idz_local;
int group_id = groupid_x + group_y*groupid_y + groupid_z*group_z*groupid_z;

int group_count = group_x*group_y*group_z;

int offset = 1;

bool isValid = true;

if(idx >sizex) isValid = false;
if(idy >sizey) isValid = false;
if(idz >sizez) isValid = false;
if(idx_local >local_x) isValid = false;
if(idy_local >local_y) isValid = false;
if(idz_local >local_z) isValid = false;

if(isValid)
{
barrier(CLK_LOCAL_MEM_FENCE);
__local float Minvals[WORKGROUPSIZE];
__local float Maxvals[WORKGROUPSIZE];

Minvals[gid_local] = input[gid];
Maxvals[gid_local] = input[gid];

int offset = 1;
int jump = 2;
barrier(CLK_LOCAL_MEM_FENCE);
for(;gid_local%jump==0;)
{
Minvals[gid_local] = fmin(Minvals[gid_local],Minvals[gid_local+offset]);
Maxvals[gid_local] = fmax(Maxvals[gid_local],Maxvals[gid_local+offset]);

offset *= 2;
jump *= 2;
}

barrier(CLK_GLOBAL_MEM_FENCE);

if(gid_local == 0)
{
//global Rueckschreiben
output[gid*2] = Minvals[gid_local];
output[gid*2 +1] = Maxvals[gid_local];

}

barrier(CLK_GLOBAL_MEM_FENCE);
if(gid==0)
{
float mint;
float maxt;
for(int i =0; i < group_count; i=i+2)
{
mint = fmin(mint,output[i]);
maxt = fmax(maxt,output[i+1]);

}

output[0] = mint;
output[1] = maxt;
}

}
}

Es gibt nen Bluescreen, welcher auf einen Konflikt mit einem Treibereset hindeutet(BCCode: 116).

Problemsignatur:
Problemereignisname: BlueScreen
Betriebsystemversion: 6.1.7601.2.1.0.256.48
Gebietsschema-ID: 1031

Zusatzinformationen zum Problem:
BCCode: 116
BCP1: FFFFFA80043354E0
BCP2: FFFFF88006E07408
BCP3: 0000000000000000
BCP4: 0000000000000002
OS Version: 6_1_7601
Service Pack: 1_0
Product: 256_1

Ich hab keine Ahnung, was das Problem ist. Bitte anderen Thread beachten.
Ob es ein Hardware oder Softwareproblem ist, kann ich nicht sagen, deswegen zwei Threads.

Zum Kernel:
Der sucht aus einem Array aus Floats das Max und Min heraus. Dabei wird erstmal in der Workgroup reduziert, dann Global.

[Gast]

Im Fall dass gid_local=0 ist, und jump !=0 (startet ja mit dem Wert 2) ist gid_local%jump immer 0. Nun multiplizierst du jump immer wieder mit 2, irgendwann gibt es da einen Überlauf und jump ist 0, nun teilst du durch 0, und ich habe im Moment keine Ahnung wie eine Grafikkarte darauf reagiert, aber könnte mir durchaus vorstellen, dass es ein Bluescreen geben könnte...

[Nakai]

Danke für den Tipp!
Ich muss evtl dann die Sprungbedingung ändern bzw. die Abbruchbedingung der Schleife ändern.

Das wäre womöglich ein Grund, warum alle meine Kernels nicht liefen, da ich immer die Abbruchbedingung mit % errechnet habe.
Gut, werd ich morgen mal wieder anschauen.

[Nakai]

Ich hab schnell einen einfache uns simplen 1D-Kernel erstellt.

Das hat geklappt. Danke, für die Info. Anscheinend klappts jetzt.

[Nakai]

Hat jemand Erfahrungen mit CodeXL?

Debugging ist unzuverlässig. Variabeln werden mit falschen Werten initialisiert.


bool isValid = true;


Wird mit dem Wert 64(!) initialisiert. Natürlich funktionieren, dann keine if-Abfragen damit. Eigentlich sollte diese Variable private sein, ergo nur für das derzeitige WorkItem bzw. für alle Workitems private.
Wenn ich dann zwischen den Workitems switche, ändert sich der Wert der Variable.

Liegt es am CodeXl oder bin ich einfach zu deppert?

[samm]

Vorneweg, ich habe nie mit OpenCL gearbeitet. Dennoch bisschen Input:


"true" kann prinzipiell irgendein Wert sein, der nicht 0 ist (für einfache Implementation der logischen Verknüpfungen scheints mir logisch so ;) ).
Natürlich funktionieren, dann keine if-Abfragen damit.Wirklich nicht? Annahme oder getestet?

Zudem ist bool laut google in OpenCL implementationstechnisch nicht fix definiert (keine Grösse festgelegt) und man muss bei Benutzung einiges beachten, vgl. ersten Suchmaschinentreffer: http://stackoverflow.com/questions/4441280/does-opencl-support-boolean-variables

Wenn ich dann zwischen den Workitems switche, ändert sich der Wert der Variable.
Wenn isValid, wie es klingt, eine lokale Variable ist, passt das doch, dass die pro WorkItem einen anderen Wert annehmen kann. Oder meinst du WorkItem1.isValid == 64, switch zu WorkItem2, switch zurück zu WorkItem1 und dann ist WorkItem1.isValid == 123?

[Nakai]

Wirklich nicht? Annahme oder getestet?

Es geht um diesen Block.

bool isValid = true;

if(idx >sizex) isValid = false;
if(idy >sizey) isValid = false;
if(idz >sizez) isValid = false;
if(idx_local >local_x) isValid = false;
if(idy_local >local_y) isValid = false;
if(idz_local >local_z) isValid = false;

if(isValid)
{


Hier prüfe ich ab, welcher Thread operieren darf. Laut Debugger springe ich nicht in diesen If-Block. isValid wird nicht geändert, also sollte er hineinspringen. Tut er nicht.
Langsam glaube ich, dass der CodeXL-Debugger Unsinn anzeigt
Ich vermute, dass es an der Kombination Treiber, Hardware und OpenCl RTE liegt, aber keine Ahnung.

Wenn isValid, wie es klingt, eine lokale Variable ist, passt das doch, dass die pro WorkItem einen anderen Wert annehmen kann. Oder meinst du WorkItem1.isValid == 64, switch zu WorkItem2, switch zurück zu WorkItem1 und dann ist WorkItem1.isValid == 123?

Ja du hast recht, es passt dazu. Das Problem ist, wenn ich die Variable als Int(int isValid = 1; 1 dann true) definiere und genauso auf Gültigkeit teste(0 dann false; if(isValid == 1)), kann es sein, dass der Int mit 0 oder anderen Werten initialisiert wird. CodeXL erkennt auch, dass es sich um ein Int handelt, aber der Wert ist falsch. Die If-Bedingung wird gekonnt ignoriert und übersprungen.

[samm]

Danke für die Ausführung! Okay, das klingt dann wirklich nach fehlerhaftem Debugger resp. fehlerhafter Kompilierung für den Debug-Mode. Ein Hardware-Problem würde ich eher ausschliessen, denke mal nicht, dass sich da was ständig verrechnet ;) Treiber, OpenCL-Environment oder der Debugger, ja - kannst du jeweils eine einzelne der Komponenten ändern und erneut testen?

[Nakai]

Ich bin ein bisschen ratlos. Werde ab morgen auf einen anderen Rechner ausweichen(Betonung auf Rechner; kein Notebook). Es ist ein Witz.

AMDs Testsoftware(Teacup) läuft ohne Probleme, eigene kleine Kernels genauso. Ich hab es mittlerweile und zeitweise geschafft, dass der Debugger in die IF springt. Danach hört der Debugvorgang irgendwann einfach auf. Siehe Anhang.

Mittlerweile wieder das typische Problem, dass er nicht in die IF springt. Wird der Code im Profiler gelauncht, gibts entweder Bluescreen oder Driverreset.

Ich hab keine Ahnung woran es liegt. Naja...anderer PC ich komme...

[Gast]

Was passiert eigentlich, wenn du den Kernel auf der Cpu ausführst?

[Nakai]

Aufgrund der Kapselung von OpenCL in ITK, ist es nicht möglich den Kernel auf der GPU auszuführen. ITK will OpenCL-Kernels nur auf GPUs ausführen. Das zu ändern, wäre kein Problem, aber dann müsst ich den ITK-Code abändern und das will ich eher nicht tun. ;)

Mittlerweile steht nun sicher fest, dass mein Notebook das Problem verursacht hat. OpenCL-Code läuft korrekt auf einer anderen Maschine(Tesla C2050). Der MinMax-Kernel wurde noch abgeändert(Fehler bereinigt) und gibt die richtigen Werte zurück(auf der Tesla). Auf dem Notebook wird nur totaller Mist berechnet.

Normalerweise sollte es gehen, habe hierzu auch an anderen Stellen gefragt.

Das Problem kann man eindeutig auf AMDs Treiberpolitik im OEM-Bereich(vor allem Mobile) schieben. Wenn man schon den Treibersupport auf die OEMs abschiebt und man selber Mobilitytreiber anbietet, sollten diese wirklich auf alle Devices auch laufen. Auch Dell kann man den schwarzen Peter zuschieben. Insgesamt kamen nur drei verschiedene Treiber von Dell für mein Notebook raus. Nur der Älteste(über ein Jahr alt) ist in der Lage OpenCL auszuführen(RTE inbegriffen), wird aber nicht von AMDs CodeXL unterstützt(Treiber zu alt, wird nicht erkannt). Schade...

[Gast]

Ok dass der Mobility Treiber nicht mit der Gpu will ist sehr ärgerlich. Vermutlich wird hier Dell die Schuld Amd geben, während Amd den schwarzen Peter Dell zuschiebt.
Interessant wäre, ob der Code korrekt auf einer normalen Amd Gpu mit aktuellen Treibern läuft, oder ob er allgemein nicht mit Amd klarkommt.