Nearest Neighbours in CUDA Particles - Algorithmus, Opencl, Simulation, Physik, nächster Nachbar

Bearbeiten 2: Bitte werfen Sie einen Blick auf dieser Kreuzweg für TLDR.

Bearbeiten: Vorausgesetzt, dass die Partikel in Gitterzellen segmentiert sind 16^3 Grid), ist es eine bessere Idee, eine Arbeitsgruppe für jede Gitterzelle und so viele Arbeitseinheiten in einer Arbeitsgruppe laufen zu lassen, wie es eine maximale Anzahl von Partikeln pro Gitterzelle geben kann?

In diesem Fall könnte ich alle Partikel aus ladenbenachbarte Zellen in den lokalen Speicher und iterieren durch sie einige Eigenschaften zu berechnen. Dann könnte ich einen bestimmten Wert in jedes Partikel in der aktuellen Rasterzelle schreiben.

Wäre dieser Ansatz von Vorteil gegenüber dem Ausführen des Kerns für alle Partikel und für jedes Iterieren über (meistens die gleichen) Nachbarn?

Auch, was ist das ideale Verhältnis von number of particles/number of grid cells?

---

Ich versuche zu re-implementieren (und zu modifizieren) CUDA-Partikel für OpenCL und verwenden Sie es, um die nächsten Nachbarn für jedes Partikel abzufragen. Ich habe folgende Strukturen erstellt:

• Puffer P Halten aller Partikel "3D-Positionen (float3)

• Puffer Sp Speicherung int2 Paare von Partikel-IDs und ihre räumlichen Hashes. Sp ist nach dem Hash sortiert. (Der Hash ist nur eine einfache lineare Zuordnung von 3D zu 1D - noch keine Z-Indizierung.)

• Puffer L Speicherung int2 Paare von Anfangs- und Endpositionen bestimmter räumlicher Hashwerte im Puffer Sp. Beispiel: L[12] = (int2)(0, 50).

  • L[12].x ist der Index (in Sp) des zuerst Partikel mit räumlichem Hash 12.
  • L[12].y ist der Index (in Sp) des letzte Partikel mit räumlichem Hash 12.

Jetzt, da ich all diese Puffer habe, möchte ich alle Partikel durchlaufen P und für jedes Teilchen iteriere durch seine nächsten Nachbarn. Momentan habe ich einen Kernel, der so aussieht (Pseudocode):

__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid             = get_global_id(0);
float3 curr_particle   = P[gid];
int    processed_value = 0;

for(int x=-1; x<=1; x++)
for(int y=-1; y<=1; y++)
for(int z=-1; z<=1; z++) {

float3 neigh_position = curr_particle + (float3)(x,y,z)*GRID_CELL_SIDE;

// ugly boundary checking
if ( dot(neigh_position<0,        (float3)(1)) +
dot(neigh_position>BOUNDARY, (float3)(1))   != 0)
continue;

int neigh_hash        = spatial_hash( neigh_position );
int2 particles_range  = L[ neigh_hash ];

for(int p=particles_range.x; p<particles_range.y; p++)
processed_value += heavy_computation( P[ Sp[p].y ] );

}

Out[gid] = processed_value;
}

Das Problem mit diesem Code ist, dass es langsam ist. Ich vermute den nichtlinearen GPU-Speicherzugriff (insbesondere P[Sp[p].y] im innersten for Schleife) verursacht die Langsamkeit.

Was ich tun möchte, ist zu benutzen Z-Ordnungs-Kurve als räumlicher Hash. So konnte ich nur 1 haben for Schleife, die bei der Abfrage von Nachbarn durch einen fortlaufenden Speicherbereich iteriert. Das einzige Problem ist, dass ich nicht weiß, was die Z-Indexwerte starten und stoppen sollen.

Der heilige Gral, den ich erreichen möchte:

__kernel process_particles(float3* P, int2* Sp, int2* L, int* Out) {
size_t gid             = get_global_id(0);
float3 curr_particle   = P[gid];
int    processed_value = 0;

// How to accomplish this??
// `get_neighbors_range()` returns start and end Z-index values
// representing the start and end near neighbors cells range
int2 nearest_neighboring_cells_range = get_neighbors_range(curr_particle);
int first_particle_id = L[ nearest_neighboring_cells_range.x ].x;
int last_particle_id  = L[ nearest_neighboring_cells_range.y ].y;

for(int p=first_particle_id; p<=last_particle_id; p++) {
processed_value += heavy_computation( P[ Sp[p].y ] );
}

Out[gid] = processed_value;
}

Antworten:

-1 für die Antwort № 1

Sie sollten die Morton Code-Algorithmen genau studieren. Ericsons Echtzeit-Kollisionserkennung erklärt das sehr gut.

Ericson - Echtzeit Kollisionserkennung

Hier ist eine weitere nette Erklärung mit einigen Tests:

Morton-Kodierung / Dekodierung durch Bit-Interleaving: Implementierungen

Z-Order-Algorithmen definieren nur die Pfade derKoordinaten, in denen Sie von 2 oder 3D Koordinaten zu nur einer Ganzzahl hashen können. Obwohl der Algorithmus für jede Iteration tiefer geht, müssen Sie die Grenzwerte selbst festlegen. Normalerweise wird der Stop-Index durch ein Sentinel angegeben. Lassen Sie den Sentinel stoppen, wird Ihnen sagen, auf welcher Ebene das Partikel platziert ist. Die maximale Ebene, die Sie definieren möchten, gibt Ihnen die Anzahl der Zellen pro Dimension an. Zum Beispiel mit maximaler Stufe bei 6 haben Sie 2 ^ 6 = 64. Sie haben 64x64x64 Zellen in Ihrem System (3D). Das bedeutet auch, dass Sie ganzzahlige Koordinaten verwenden müssen. Wenn Sie Floats verwenden, müssen Sie wie konvertieren coord.x = 64*float_x und so weiter.

Wenn Sie wissen, wie viele Zellen Sie in Ihrem System haben, können Sie Ihre Grenzen definieren. Versuchen Sie einen binären Octree zu verwenden?

Da Teilchen in Bewegung sind (in diesem CUDA-Beispiel), sollten Sie versuchen, über die Anzahl der Teilchen anstelle von Zellen zu parallelisieren.

Wenn Sie Listen der nächsten Nachbarn erstellen möchtenSie müssen die Partikel Zellen zuordnen. Dies geschieht durch eine Tabelle, die anschließend von Zellen nach Partikeln sortiert wird. Trotzdem sollten Sie die Partikel durchlaufen und auf ihre Nachbarn zugreifen.

Über deinen Code:

Das Problem mit diesem Code ist, dass er langsam ist. Ich vermute, dass der nichtlineare GPU-Speicherzugriff (insbesondere P [Sp [p] .y) in der innersten for-Schleife) die Langsamkeit verursacht.

Erinnern Sie sich an Donald Knuth. Sie sollten messen, wo der Flaschenhals ist. Sie können den NVCC Profiler verwenden und nach Engpässen suchen. Nicht sicher, was OpenCL als Profiler hat.

    // ugly boundary checking
if ( dot(neigh_position<0,        (float3)(1)) +
dot(neigh_position>BOUNDARY, (float3)(1))   != 0)
continue;

Ich denke, du solltest es nicht so abzweigen, wie wäre es mit Null, wenn du anrufst heavy_computation. Nicht sicher, aber vielleicht hast du hier eine Art Verzweigungsprognose. Versuchen Sie das irgendwie zu entfernen.

Parallel über die Zellen zu laufen ist eine gute Ideenur wenn Sie keine Schreibzugriffe auf die Partikeldaten haben, ansonsten müssen Sie Atomics verwenden. Wenn Sie stattdessen über den Partikelbereich gehen, lesen Sie Zugriffe auf die Zellen und Nachbarn, aber Sie erstellen Ihre Summe parallel, und Sie sind nicht gezwungen, ein Race-Condition-Paradigma zu wählen.

Was ist das ideale Verhältnis von Anzahl der Partikel / Anzahl der Gitterzellen?

Kommt wirklich auf deine Algorithmen und die anPartikelpackung innerhalb Ihrer Domäne, aber in Ihrem Fall würde ich die Zellgröße definieren, die dem Partikeldurchmesser entspricht, und nur die Anzahl der Zellen verwenden, die Sie erhalten.

Wenn Sie also die Z-Ordnung verwenden und Ihren heiligen Gral erreichen möchten, versuchen Sie, ganzzahlige Koordinaten zu verwenden und diese zu hashen.

Versuchen Sie auch, größere Mengen an Partikeln zu verwenden. Etwa 65000 Partikel, wie sie in CUDA-Beispielen verwendet werden, sollten Sie berücksichtigen, da auf diese Weise die Parallelisierung größtenteils effizient ist. Die laufenden Verarbeitungseinheiten werden ausgenutzt (weniger Leerlauf-Threads).

---