Sparse Array in CUDA oder OpenCL

Question

Ich habe ein großes Array (z.B. 512K Elemente), das auf der GPU residiert und bei dem nur ein kleiner Teil der Elemente (z.B. 5K zufällig verteilte Elemente - Menge S) verarbeitet werden muss. Der Algorithmus, um herauszufinden, welche Elemente zu S gehören, ist sehr effizient, so dass ich leicht ein Array A mit Zeigern oder Indizes auf Elemente aus der Menge S erstellen kann.

Was ist der effizienteste Weg, um einen CUDA- oder OpenCL-Kernel nur über Elemente von S auszuführen? Kann ich einen Kernel über ein Array A ausführen? Alle Beispiele, die ich bisher gesehen habe, befassen sich mit zusammenhängenden 1D-, 2D- oder 3D-Arrays. Gibt es ein Problem mit der Einführung einer indirekten Ebene?

Answer 1

In CUDA wird der zusammenhängende (nicht zufällige) Speicherzugriff aufgrund der möglichen Verwendung von Memory Coalescing bevorzugt. Es ist keine große Sache, ein Array von zufällig verteilten Indizes zu erstellen und einen Index von A pro Thread zu bearbeiten, etwa so:

__global__ kernel_func(unsigned * A, float * S)
{
    const unsigned idx = threadIdx.x + blockIdx.x * blockDim.x;
    const unsigned S_idx = A[idx];

    S[S_idx] *= 5; // for example...
    ...
}

Der Speicherzugriff auf S[random access] wird jedoch sehr langsam sein (hier wird ein möglicher Engpass liegen).

Wenn Sie sich für CUDA entscheiden, dann müssen Sie viel mit Block-/Gittergrößen experimentieren, den Registerverbrauch pro Thread minimieren (um die Anzahl der Blöcke pro Multiprozessor zu maximieren) und vielleicht A sortieren, um das nächstgelegene S_ind von den nächstgelegenen Threads zu verwenden...

Answer 2

Wenn Sie Ihre Indizes sortieren oder die Liste sortiert bauen, die Leistung viel helfen wird, wenn es Cluster von Indizes dann versuchen, mit Textur-Speicher, und wenn Sie Zugriff auf eine Reihe von Elementen aus jedem Thread mit einigen über Runde die ich gefunden mit dem gemeinsamen Speicher gibt eine erhebliche Leistungssteigerung.

Answer 3

Die eine Umleitungsebene ist überhaupt kein Problem. Ich verwende das ziemlich oft in meinem eigenen CUDA Code. Wird die Menge S im Laufe der Zeit wahrscheinlich statisch bleiben? Wenn ja, kann es sich durchaus lohnen, das Lookup A zu generieren, wie Sie sagten.

Außerdem ist der Texturspeicher Ihr Freund bei der Bereitstellung von Cache-Lokalität. Die Art der Textur, die Sie verwenden (1D, 2D oder 3D), hängt von Ihrem Problem ab.