OpenCL: Optimale Work-Group Size

[Nakai]

Heyho,
ich hätte mal eine Frage bezüglich der optimalen Work-Group Size bei OpenCL. Die Größer einer Work-Group ist ja architekturspezifisch auf die Hardware ausgelegt. Bzw. die maximale Größe hängt von der Hardware ab. Wie entscheidet man am besten, welche Größe die Workgroups haben sollten?
Umso mehr lokale Werte und private Variabeln ein WorkItem(Workgroup) hat, desto kleiner sollten diese sein.

Derzeit verwende ich eine Work-Group Size von 64(3D: 4 in jede Richtung), 256(2D:16 je DIM) und 256(1D). Das ist natürlich komplett unabhängig und wird per #define als Preambel gesetzt.

Gibt es irgendwelche Richtlinien auf was man beim Festlegen der Work-Group Size beachten muss?

[Gast]

Heyho,
ich hätte mal eine Frage bezüglich der optimalen Work-Group Size bei OpenCL. Die Größer einer Work-Group ist ja architekturspezifisch auf die Hardware ausgelegt. Bzw. die maximale Größe hängt von der Hardware ab. Wie entscheidet man am besten, welche Größe die Workgroups haben sollten?
Umso mehr lokale Werte und private Variabeln ein WorkItem(Workgroup) hat, desto kleiner sollten diese sein.
das hat nicht viel mit 'sollte' zu tun, sondern mit hardware limits. z.b. gibt es ein limit an registern, entsprechend hast du
work-group-size = HardwareReigster/KernelRegister;


Derzeit verwende ich eine Work-Group Size von 64(3D: 4 in jede Richtung), 256(2D:16 je DIM) und 256(1D). Das ist natürlich komplett unabhängig und wird per #define als Preambel gesetzt. unabhaengig und define?
1. du solltest abfragen wie gross eine work group sein kann, das bekommst du per kernel von opencl. gibst du mehr an bekommst du einen fehler.
2. CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE beachten
3. mach das flexibel und nicht per define, das ist fuer einen test ok, aber nicht bei einem tool das du nutzt, willst du es neubauen wenn jemand es mal auf anderer hardware testen will?


Gibt es irgendwelche Richtlinien auf was man beim Festlegen der Work-Group Size beachten muss?
ja, die ignoriere theoretische richtlinien, der einzige weg das wirklich zu wissen ist es zu profilen. manchmal kann eine kleinere workgroud size schneller sein, weil du mit grossen cache trasht, aber sobald deine daten in den cache passen, macht verkleinern vielleicht wieder alles ineffizienter, weil weniger threads zum latenzverstecken vorhanden sind....

oft ist es sogar von den daten abhaengig, kann also sein dass der kernel mit den einen daten eine andere workgroup size bevorzugt als mit anderen daten.

ich habe immer ein kleines testbed im program, beim start lauf ich damit verschiedene tests,
1. genauigkeit (je nach implementierung kann es sehr abweichen, manchmal gibt es auch einfach bugs und du willst nicht bugreports usw. um dann nach ewigkeiten festzustellen die fremde hw bzw driver war kaputt).
2. performance je nach item groesse
3. performance nach device (gerade mit den samples bekommt man den eindruck immer das device waehlen zu sollen das die meisten compute units hat etc. aber du wirst oft feststellen, dass die CPU schneller sein kann als die igps), ein kleiner benchmark der die performance testet damit man das device waehlt was am schnellsten ist kann echt viel bringen.

[Nasenbaer]

Dem kann ich zustimmen. Hatte auch schon mal Optimierungen an einem Kernel vorgenommen nur um dann festzustellen, dass es danach langsamer läuft. Hatte einfach die Fähigkeiten des Caches unterschätzt.
Falls du AMD-H/W nutzt, dann kann ich dir den AMD APP KernelAnalyzer, APP Kernel Profiler und ggf. CodeXL (derzeit nicht Win8 x64) empfehlen. Mit NVIDIA ist man AFAIK bei OpenCL aufgeschmissen. Deren NSight ist ein super Programm aber kann wohl nur mit CUDA umgehen bzw. bei OpenCL nur ganz wenig anzeigen.

[Nakai]

Danke, für die Antworten.

unabhaengig und define?
1. du solltest abfragen wie gross eine work group sein kann, das bekommst du per kernel von opencl. gibst du mehr an bekommst du einen fehler.
2. CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE beachten
3. mach das flexibel und nicht per define, das ist fuer einen test ok, aber nicht bei einem tool das du nutzt, willst du es neubauen wenn jemand es mal auf anderer hardware testen will?

Das meinte ich so. Ich erstelle in gewisse Kernels Local Memory mit einer bestimmten Größe. Diese Größe entspricht der Größe der WorkgroupSize. Ergo pro WorkItem ein Platz im Array. Je nach Problem wird das natürlich angepasst.
Es ist prinzipiel flexibel gedacht, da ich vor dem Kompilieren des Kernels die optimale WorkgroupSize bestimmen möchte. Aber anscheinend ist das nicht so einfach.

ja, die ignoriere theoretische richtlinien, der einzige weg das wirklich zu wissen ist es zu profilen. manchmal kann eine kleinere workgroud size schneller sein, weil du mit grossen cache trasht, aber sobald deine daten in den cache passen, macht verkleinern vielleicht wieder alles ineffizienter, weil weniger threads zum latenzverstecken vorhanden sind....

Danke, sowas wollte ich hören. Dann weiß ich endlich, wie ich das Problem anpacken muss.

ich habe immer ein kleines testbed im program, beim start lauf ich damit verschiedene tests,
1. genauigkeit (je nach implementierung kann es sehr abweichen, manchmal gibt es auch einfach bugs und du willst nicht bugreports usw. um dann nach ewigkeiten festzustellen die fremde hw bzw driver war kaputt).
2. performance je nach item groesse
3. performance nach device (gerade mit den samples bekommt man den eindruck immer das device waehlen zu sollen das die meisten compute units hat etc. aber du wirst oft feststellen, dass die CPU schneller sein kann als die igps), ein kleiner benchmark der die performance testet damit man das device waehlt was am schnellsten ist kann echt viel bringen.

=)

Das ist ein guter Hinweis, welchen ich befolgen werde.

Dem kann ich zustimmen. Hatte auch schon mal Optimierungen an einem Kernel vorgenommen nur um dann festzustellen, dass es danach langsamer läuft. Hatte einfach die Fähigkeiten des Caches unterschätzt.
Falls du AMD-H/W nutzt, dann kann ich dir den AMD APP KernelAnalyzer, APP Kernel Profiler und ggf. CodeXL (derzeit nicht Win8 x64) empfehlen. Mit NVIDIA ist man AFAIK bei OpenCL aufgeschmissen. Deren NSight ist ein super Programm aber kann wohl nur mit CUDA umgehen bzw. bei OpenCL nur ganz wenig anzeigen.

Ich bin von Nvidias Engagement bei OpenCL wirklich angepisst. Klar, man will das eigene CUDA pushen. Da hat AMD deutlich bessere Arbeit geleistet.

[del_4901]

Ich bin von Nvidias Engagement bei OpenCL wirklich angepisst. Klar, man will das eigene CUDA pushen. Da hat AMD deutlich bessere Arbeit geleistet.OpenCL war einfach eine Totgeburt, das Versagen hat sich die Kronos Gruppe zuzuschreiben. Selbst AMD entwickelt mit HSA die bessere Alternative.

[Coda]

HSA ist doch nur eine Architektur, was hat das mit der Programmierumgebung zu tun?

[del_4901]

HSA ist doch nur eine Architektur, was hat das mit der Programmierumgebung zu tun?Die Architektur ist nur der Antrieb, bei HSA handelt es sich aber um ein Komplettpaket.
http://developer.amd.com/resources/heterogeneous-computing/what-is-heterogeneous-system-architecture-hsa/

[Coda]

Exakt. Es ist ein Komplettpaket, dass unter anderem auch OpenCL enthält:
AMD is starting this process by delivering HSA optimized programming tools for today’s most widely available heterogeneous languages: OpenCL™ and C++ AMP.

[del_4901]

Exakt. Es ist ein Komplettpaket, dass unter anderem auch OpenCL enthält: Nur weil es enthalten ist, macht es OpenCL noch lange nicht besser. Du kannst bei HSA ein beliebiges Frontend reindruecken, und OpenCL waehre dabei nicht meine erste Wahl.

The HSA team at AMD analyzed the performance of Haar Face Detect, a commonly used multi-stage video analysis algorithm used to identify faces in a video stream. The team compared a CPU/GPU implementation in OpenCL™ against an HSA implementation. The HSA version seamlessly shares data between CPU and GPU, without memory copies or cache flushes because it assigns each part of the workload to the most appropriate processor with minimal dispatch overhead. The net result was a 2.3x relative performance gain at a 2.4x reduced power level*. This level of performance is not possible using only multicore CPU, only GPU, or even combined CPU and GPU with today’s driver model. Just as important, it is done using simple extensions to C++, not a totally different programming model. Und such dir nicht immer die "Brocken" aus dem Artikel raus die gerade irgendwie in deine Argumentationkette passen. Man kann auch ruhig mal sagen, ok war mir nicht bewust. Sowas kommt bei mir wesentlich kompetenter an.

[Nakai]

Dumme Frage:
Ich versuche gerade einen LowPassFilter zu parallelisieren. Die sequentielle Berechnung auf der CPU liefert teilweise deutlich andere Ergebnisse, teilweise die gleichen oder sehr ähnliche Ergebnisse. Ich vermute(falls mein LowpassAlgo nicht völlig falsch ist; er filtert tatsächlich), dass OpenCL nicht wie die CPU "rundet". Achja der CPU-Code hat keinerlei Optimierungen(wurde nicht von mir geschrieben), also bestimmte Flags in der Makefile werden nicht gesetzt(kein -mfpmath=sse, o.ä). Es läuft alles anscheinend auf X87.

Die Unterschiede zwischen CPU und GPU-Ergebnisse sind teilweise unfassbar gleich, andererseits weichen sie ziemlich stark ab. Ich werde meinen Code noch auf Fehler überprüfen, jedoch gehe ich mitterweile stark davon aus, dass die GPU-Implementierung falsch rundet. Gibt es Möglichkeiten das zu unterbinden?


mfg

[RLZ]

und OpenCL waehre dabei nicht meine erste Wahl.
Was wäre deine erste Wahl?

[del_4901]

Was wäre deine erste Wahl?
Haskell :ugly:
ok, doch lieber, OCaml.
Man muesste nur ein eigenes Frontend/Debugger Toolchain fuer die HSA-IL schreiben.
Aber selbst DX Compute oder C++ AMP ist besser als OpenCL

Die Unterschiede zwischen CPU und GPU-Ergebnisse sind teilweise unfassbar gleich, andererseits weichen sie ziemlich stark ab. Ich werde meinen Code noch auf Fehler überprüfen, jedoch gehe ich mitterweile stark davon aus, dass die GPU-Implementierung falsch rundet. Gibt es Möglichkeiten das zu unterbinden? k.A. aber ich denke nicht das man die Rundung umschalten kann. Was rundest du denn f32? Es hoert sich fuer mich eher nach einer Racecondition an.

[del_4901]

Haskell :ugly:
Wo bleibt der Shitstorm? Ich hatte mich schon so drauf gefreut. ;)

[Nakai]

k.A. aber ich denke nicht das man die Rundung umschalten kann. Was rundest du denn f32? Es hoert sich fuer mich eher nach einer Racecondition an.

Es sind Rundungsfehler. Das Problem ist mein Referenzcode, welcher keine Erweiterungen benutzt(keine Compilerflags für SSE etc.). Der läuft auf dem klassischen X87. Anscheinend verwendet X87 intern immer die 80Bit-Präzision. Die GPU hat nur 32Bit bzw. 64Bit Register(bei DP). Da zerlegt es jegliche Operation auf der GPU. Das Problem auf dem X87 ist jedoch das Rückspeichern in normale Floats. Da wird abgeschnitten bzw. rückkonvertiert. Dennoch sollte das Ergebnis auf der CPU genauer sein, einfach weil viel mehr Bits zur Verfügung sind.

Das ist hart, hätte ich nicht gedacht. :eek:

http://docs.oracle.com/cd/E19957-01/806-3568/ncg_goldberg.html

http://www.securistreet.com/dri/up/2_FPU.pdf

Bei iterative Verfahren summieren sich wohl dadurch die Abweichungen und die Fehler bzw. können sich auch auslöschen. Außerdem ist es verdammt wichtig in welcher Reihenfolge die Operationen stattfinden.

[del_4901]

Bei iterative Verfahren summieren sich wohl dadurch die Abweichungen und die Fehler bzw. können sich auch auslöschen. Außerdem ist es verdammt wichtig in welcher Reihenfolge die Operationen stattfinden. Das zum einen, das kann man aber loesen indem man nicht gerade kleine mit grossen Zahlen zusammenrechnet. Zum anderen kann es auch einfach sein, dass du dir NaNs oder Denorms einfaengst. Dann bekommst du undefiniertes Verhalten. Mach mal beim x87 code floatingpoint exceptions an, vllt. ist der schon "kaputt", und man merkt es nur nicht. Und danach kannst du ja mal die Präzision der x87 FPU verstellen, wenn dann der Fehler nicht auftritt, dann wuerde ich den Fehler eher in deinem Code vermuten.

[Nakai]

Das zum einen, das kann man aber loesen indem man nicht gerade kleine mit grossen Zahlen zusammenrechnet. Zum anderen kann es auch einfach sein, dass du dir NaNs oder Denorms einfaengst. Dann bekommst du undefiniertes Verhalten. Mach mal beim x87 code floatingpoint exceptions an, vllt. ist der schon "kaputt", und man merkt es nur nicht.

Mhh, ich glaube nicht, dass der X87-Code kaputt ist. Es gibt im Code schon einige Optimierungsverfahren um die Zahlen zu optimieren.

http://www.fmrib.ox.ac.uk/analysis/research/fast

Das mit den NaNs und Denorms ist auch unwahrscheinlich. Das wird abgeprüft.

Das mit kleine und große Zahlen wird auch ein Grund sein. Die Hauptschuld gebe ich dem iterativen Verfahren.

Und danach kannst du ja mal die Präzision der x87 FPU verstellen, wenn dann der Fehler nicht auftritt, dann wuerde ich den Fehler eher in deinem Code vermuten.

Das wäre interessant, aber wie? Oder meinst du, ich soll auf SSE umstellen?

[del_4901]

Mhh, ich glaube nicht, dass der X87-Code kaputt ist. Es gibt im Code schon einige Optimierungsverfahren um die Zahlen zu optimieren.

http://www.fmrib.ox.ac.uk/analysis/research/fast

Das mit den NaNs und Denorms ist auch unwahrscheinlich. Das wird abgeprüft.

Das mit kleine und große Zahlen wird auch ein Grund sein. Die Hauptschuld gebe ich dem iterativen Verfahren. Glauben kannst du in der Kirche! Du sollst aber messen! Es waehre ja nicht das erste mal das eine Lib nicht ganz rund laueft.



Das wäre interessant, aber wie? Oder meinst du, ich soll auf SSE umstellen? Es gibt spezielle ASM instruktionen, mit denen du sowas pro thread verstellen kannst.
http://www.altdevblogaday.com/2012/04/20/exceptional-floating-point/
http://www.altdevblogaday.com/2012/03/22/intermediate-floating-point-precision/

[Nakai]

Glauben kannst du in der Kirche! Du sollst aber messen! Es waehre ja nicht das erste mal das eine Lib nicht ganz rund laueft.

Das wäre übel, da diese Lib tatsächlich in der Medizin verwendet wird. Ich kann mir das nur seeehr schwer vorstellen. Und ja, ich werde das nachprüfen.

Es gibt spezielle ASM instruktionen, mit denen du sowas pro thread verstellen kannst.

Also doch per Assembler. Werd ich mir morgen zur Gemüte führen. Danke, für die Links.

[Gast]

Ich hatte noch nie Probleme mit der Präzision GPU vs. CPU, und wir haben fast ausschließlich iterative Algorithmen. Noch dazu auf der GPU alles in float! Der Fehler zu CPU double lag auch bei >500 Iterationen max im Bereich 1e-4 typischerweise eher 1e-6 und kleiner, also gerade die Präzision die man mit float überhaupt erreichen kann. Sprich das Ergebnis war identisch.

Mal gleiches Ergebnis und mal große (!) Abweichung hört sich für mich nach einem typischen Fehler im parallelen Code an. Concurrent writes/reads ohne Synchronisation und die Threads lesen halt irgendwelche falschen Daten oder sowas.

[Nakai]

Ich habe gerade den X87-Modus angeguckt. Die läuft in 80Bit. Außerdem wird keinerlei SSE verwendet. Es wird kein Flag in der Makefile gesetzt.

Ich hatte noch nie Probleme mit der Präzision GPU vs. CPU, und wir haben fast ausschließlich iterative Algorithmen. Noch dazu auf der GPU alles in float! Der Fehler zu CPU double lag auch bei >500 Iterationen max im Bereich 1e-4 typischerweise eher 1e-6 und kleiner, also gerade die Präzision die man mit float überhaupt erreichen kann. Sprich das Ergebnis war identisch.

Die Ergebnisse sind praktisch identisch...bei einem Durchlauf. Da sind die Abweichungen tatsächlich eher 1e-6. Es handelt sich bei meinem OpenCL-Code um einen LowpassFilter, in ZXY. Jeder Kernel wird unabhängig ausgeführt. Welche Voxel gefiltert werden, wird anhand des ursprünglichen Bildes festgelegt. 0-Werte werden nicht gefiltert. Ansonsten wird für jede Richtung ein LocalMemory angelegt. Kurz pro Säule(in XYZ) werden die Werte hineingespeichert. Dort kann ich synchronisieren, was ich auch tue.

Mal gleiches Ergebnis und mal große (!) Abweichung hört sich für mich nach einem typischen Fehler im parallelen Code an. Concurrent writes/reads ohne Synchronisation und die Threads lesen halt irgendwelche falschen Daten oder sowas.

Mhh, die großen Abweichungen habe ich tatsächlich bekommen, wenn die Synchronisation falsch lief. Aber dann waren die Ergebnisse völlig falsch.

Ansonsten, ein Code sagt mehr als tausend Worte:

Für Z:

__kernel void LowpassZ(
__global float * input,
__global float * mri,
int size_z,
int distance,
int size,
int lowz
)
{
int idx = get_global_id(0);
int idy = get_global_id(1);
int idz = get_global_id(2);

int sizex = get_global_size(0);
int sizey = get_global_size(1);
int sizez = get_global_size(2);

int local_id = get_local_id(2);
int gid = idx+sizex*idy+sizex*sizey*local_id;

bool HasNext = false;
bool HasPre = false;
bool isValid = false;

__local float workarray[WORKGROUPSIZE];
float val = 0.0;

workarray[local_id] = 0.0; //local_mem set undirty
//am I in the image?
if(gid < size && idz < size_z)
{
if(!mri[gid])
{
isValid = false;
}
else
{
isValid = true;
workarray[local_id] = input[gid];
}
}
else
{
isValid = false;
}

if(idz > 0)
{
HasPre = (mri[gid-distance]!=0);
}
else
{
HasPre = false;
}

if(idz < size_z-1)
{
HasNext = (mri[gid+distance]!=0);
}
else
{
HasNext = false;
}



for(int iter = 0; iter<lowz; iter++)
{
float OwnValue = workarray[local_id];
float NextValue = workarray[local_id+1];
float PreValue = workarray[local_id-1];

if(isValid)
{
if(idz == 0)
{
if(HasNext)
{
OwnValue = (NextValue + OwnValue + OwnValue)/3;
}
}
else if(idz == size_z-1)
{
if(HasPre)
{
OwnValue = (PreValue + OwnValue + OwnValue)/3;
}
}
else if(idz > 0 && idz < size_z-1)
{
if(HasPre)
{
if(HasNext)
{
OwnValue = (PreValue + OwnValue + OwnValue + NextValue)/4;
}
else
{
OwnValue = (OwnValue + OwnValue + PreValue)/3;
}
}
else if(HasNext)
{
OwnValue = (OwnValue + OwnValue + NextValue)/3;
}
}
}
//Writeback
barrier(CLK_LOCAL_MEM_FENCE|CLK_GLOBAL_MEM_FENCE);
workarray[local_id] = OwnValue;
}

barrier(CLK_LOCAL_MEM_FENCE);
//Write back to global memory
if(gid < size && idz < size_z)
input[gid] = workarray[local_id];
//local_iter = input[gidz];
}
Für X:
__kernel void LowpassX(
__global float * input,
__global float * mri,
int distance,
int size_x,
int size
)
{
int sizex = get_global_size(0); //Länge LOCALSIZE
int sizey = get_global_size(1);
int sizez = get_global_size(2);

int local_id = get_local_id(0);

int idx = get_global_id(0);
int idy = get_global_id(1);
int idz = get_global_id(2);

int gid = local_id + idy*size_x + idz*size_x*sizey;

__local float workarray[WORKGROUPSIZE];

workarray[local_id] = input[gid];

barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

bool isValid = false;
bool HasNext = false;
bool HasPre = false;

if(gid < size && idx < size_x)
{
if(!mri[gid])
{
isValid = false;
}
else
{
isValid = true;
workarray[local_id] = input[gid];
}
}
else
{
isValid = false;
}

if(idx > 0)
{
HasPre = (mri[gid-distance]!=0);
}
else
{
HasPre = false;
}

if(idx < size_x-1)
{
HasNext = (mri[gid+distance]!=0);
}
else
{
HasNext = false;
}

float OwnValue = workarray[local_id];
float NextValue = workarray[local_id+1];
float PreValue = workarray[local_id-1];

barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
if(isValid)
{
if(idx == 0)
{
if(HasNext)
{
OwnValue = (NextValue + OwnValue + OwnValue)/3;
}
}
else if(idx == size_x-1)
{
if(HasPre)
{
OwnValue = (PreValue + OwnValue + OwnValue)/3;
}
}
else if(idx > 0 && idx < size_x-1)
{
if(HasPre)
{
if(HasNext)
{
OwnValue = (PreValue + OwnValue + OwnValue + NextValue)/4;
}
else
{
OwnValue = (OwnValue + OwnValue + PreValue)/3;
}
}
else if(HasNext)
{
OwnValue = (OwnValue + OwnValue + NextValue)/3;
}
}
}

barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
workarray[local_id] = OwnValue;
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

if(idx < size_x)
input[gid] = workarray[local_id];

float Checkval = input[gid];

float blub = Checkval;
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
}
Für Y:
__kernel void LowpassY(
__global float * input,
__global float * mri,
int distance,
int size_y,
int size
)
{
int sizex = get_global_size(0);
int sizey = get_global_size(1); //Länge LOCALSIZE
int sizez = get_global_size(2);

int local_id = get_local_id(1);

int idx = get_global_id(0);
int idy = get_global_id(1);
int idz = get_global_id(2);

int gid = idx + local_id*sizex + idz*size_y*sizex;

__local float workarray[WORKGROUPSIZE];

workarray[local_id] = 0;//input[gid];

barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

bool isValid = false;
bool HasNext = false;
bool HasPre = false;

float mri_val = mri[gid];

if((gid < size) && (idy < size_y))
{
if(mri_val != 0)
{
isValid = true;
workarray[local_id] = input[gid];
}
else
{
isValid = false;
}
}
else
{
isValid = false;
}

if(idy > 0 && idy< size_y)
{
HasPre = (mri[gid-distance]!=0);
}
else
{
HasPre = false;
}

if(idy >= 0 && idy < (size_y-1))
{
HasNext = (mri[gid+distance]!=0);
}
else
{
HasNext = false;
}


float OwnValue = workarray[local_id];
float NextValue = workarray[local_id+1];
float PreValue = workarray[local_id-1];


barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

if(isValid)
{
//OwnValue = workarray[local_id];
if(idy == 0) //first
{
if(HasNext)
{
OwnValue = (NextValue + OwnValue + OwnValue)/3;
}
}
else if(idy == size_y-1) //last
{
if(HasPre)
{
OwnValue = (PreValue + OwnValue + OwnValue)/3;
}
}
else if(idy > 0 && idy < size_y-1) //between
{
if(HasPre)
{
if(HasNext)
{
OwnValue = (PreValue + OwnValue + OwnValue + NextValue)/4;
}
else
{
OwnValue = (OwnValue + OwnValue + PreValue)/3;
}
}
else if(HasNext)
{
OwnValue = (OwnValue + OwnValue + NextValue)/3;
}
}
}

barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
workarray[local_id] = OwnValue;
barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

if(idy < size_y && gid < size)
input[gid] = workarray[local_id];

float Checkval = input[gid];

float blub = Checkval;

}



Ein paar Worte zum Code. Fürs Debugging ist er noch etwas aufgebläht.
Alle Arrays sind als 1D-Arrays abgespeichert(Speedup). Es wird immer geschaut, welche Voxel 0 sind. Dies geschieht Anhand des Originalbildes. Der eigentliche Tiefpassfilter wird auf ein anderes Array anewendet.
Dementsprechend wird gefiltert. D.h. dass ein Voxel mit 2 Nachbarn(bei denen die Werte im Originalbild(je nach XYZ) nicht 0 sind), wird durch 4 geteilt. Der eigene Wert wird doppelt genommen. Hat ein Voxel nur einen Nachbarn, wird nur durch 3 geteilt. Mit AMD CodeXL wurde wirklich debuggt und parallel mit dem Originalcode verglichen. Die Werte sind wirklich identisch(aber die Abweichung von 10e-4...6). Der Fehler wird auch im Zentrum des Bildes immer stärker, weil dort die Fehler von den umliegenden Voxeln deutlich stärker beeinflussen. Pro Iteration(in Z wird das Anhand des Pixelradius getan; XY laufen 100mal durch; außerhalb des Kernels) wirken sich die Werte ein Voxel weiter aus. Nach 100 Iterationen wird also 100 Voxel in XY gefiltert.

[Nakai]

So, ich hab jetzt bei der Referenzsoftware die Genauigkeit gesenkt. Also statt die typischen 80bit auf 32bit. Die Ergebnisse sind dadurch anders, wie auch nicht anders zu erwarten war. Mhh, was ist nun die richtigen Ergebnisse. Läuft es mit X87 mit 80Bit Präzision, entstehen Fehler. Das kann ich mit einer GPU nicht nachbilden...

Ansonsten habe ich die Kernels abgeändert. Ich schreibe nicht mehr in den gleichen Buffer zurück, sondern in einen extra Buffer. Die Abweichung sind nun verdammt minimal geworden. Ich weiß nicht, was da passiert ist, aber ich bin ziemlich froh, dass es endlich gute annehmbare Ergebnisse sind. :D

Die Abweichungen sind dennoch da, aber in einem realtiv geringen Bereich(~10e-3).

[del_4901]

Ansonsten habe ich die Kernels abgeändert. Ich schreibe nicht mehr in den gleichen Buffer zurück, sondern in einen extra Buffer. Die Abweichung sind nun verdammt minimal geworden. Ich weiß nicht, was da passiert ist, aber ich bin ziemlich froh, dass es endlich gute annehmbare Ergebnisse sind. :D
Ich hab dir schon gesagt was da passiert ist, das ist ne ganz klassische Racecondition die du dir damit gebaut hast.

[Nakai]

Das hätte ich aber wirklich nicht erwartet, da ich immer im Local Memory arbeite, welchen ich synchronisieren kann. Außerdem greift jedes Workitem auf seinen eigenen Speicher zu.

Und ja du hattest anscheinend recht. :)