OpenCL: Welche Art von Speicher zu verwenden?

Ich habe eine Art von Filterkern, etwa wie folgt:OpenCL: Welche Art von Speicher zu verwenden?

__kernel void filterKernel (__global float4 *filter, __global float4* in_array, __global float4* out_array) 
{ 
... 
out_array[tid] = in_array[tid] * filter[fid]; 
... 
}

Kernel filterKernel heißt mehrmals (etwa 1000-mal).
Variable Filter ist ein Array von Schwimmern, dass nie ändert es Werte (bleibt gleich für alle Arbeitsgruppen und für alle Kernel-Anrufe).
In_array enthält 32768 Schwimmer.

Was ist die beste Erklärung dieser Variablen Filter? __Konstante? __lokal? Vielleicht ein "const" hier und dort platzieren? Was hilft dem Compiler am meisten? Was macht den Code am schnellsten?

Quelle

2011-01-07 Frizz

Wurde Ihre Frage beantwortet? Wenn ja, dann akzeptiere bitte eine Antwort. – koan

Sie sollten den konstanten Adressraum (__constant) verwenden, da die meisten GPUs spezielle Caches für konstanten Speicher haben. Das einzige Problem ist, dass der konstante Speicher klein ist (in der Größenordnung von 16-64KBs).

Quelle

2011-01-07 14:13:03

Wenn ich konstanten Speicher verwende: Sollte ich eine Definition (zB ___constant float4 * filter = {1.0, 2.0, 3.0, ...};) in meiner .cl Datei hinzufügen, oder * filter * in meiner Funktionsschnittstelle behalten und füllen das Array auf der CPU bevor ich die Funktion anrufe? Gibt es überhaupt einen Unterschied? – Frizz

@Frizz: Füllen Sie es mit clSetKernelArg von der CPU. –

Das ist viel langsamer als das Füllen des Array __constant in der .cl-Datei. Siehe meine Kommentare unten. – Frizz

Wenn es nicht zu groß ist, versuchen Sie, Ihren Filter global in der .cl-Datei zu definieren.
Dort können Sie versuchen, es entweder in der __constant oder __local Platz zuordnen und vergleichen, welche schneller ist. Aber nicht alle SDKs unterstützen globale Variablen im __local Adressraum (ich schaue dir ATI an).

Wenn Sie den Filter immer noch als Kernelargument übergeben möchten, sollten Sie erwägen, dessen SetKernelArg (0, ...) nur einmal aufzurufen. Es ist nicht notwendig, SetKernelArg() 1000 mal aufzurufen, solange sich der Wert oder der Index des Kernelarguments nicht ändert. Obwohl dies keinen messbaren Einfluss auf die Leistung hat, ist es dennoch sauberer.

Quelle

2011-01-07 17:34:38 vobject

Von all den verschiedenen Methoden, die ich bisher probiert habe (konstanter Speicher, LDS, Textur-Cache), definiere ich es in der .cl-Datei als __ konstant, bei weitem (!) Am schnellsten. Überraschenderweise gibt es einen großen (!) Unterschied zwischen A), der ein __konstantes Array über SetkernelArg() an den Kernel schickt, bevor er es aufruft - oder B), das __constant Array in der .cl-Datei selbst definiert. Leider funktioniert Methode B) nur bei ATI/AMD. Nvidias OpenCL-Implementierung ist hier fehlerhaft (ich habe diesen Fehler bereits gemeldet, wird in CUDA 3.3 behoben). – Frizz

Ja, ich bin das zweite. Das CUDA SDK ist sehr fehlerhaft. Ich habe versucht, eine Byte_swap() - Funktion innerhalb des Kernels mit Bit-Shifting-Operationen zu implementieren. Der OpenCL-Compiler erkennt den Zweck der Funktion und versucht, sie auf einen bswap32 resp. bswap64-Anweisung - die nicht gefunden oder aufgelöst werden kann ... – vobject

Frizz: Haben Sie die Auswirkung der Einstellung des Kernel-Arguments für den __konstanten Puffer nur einmal im Vergleich zu vor jedem Kernel-Enqueue überprüft? – dietr

__local wäre falsch, da Sie es zu nichts initialisieren können. Sie möchten wahrscheinlich __constant verwenden, vorausgesetzt, dass es passt.

Quelle

2011-01-07 20:55:38 arsenm

OpenCL: Welche Art von Speicher zu verwenden?

Antwort

Verwandte Themen