Wie effizient Daten von Threads in CUDA sammeln?

Question

Ich habe eine Anwendung, die ein System von Gleichungen in CUDA löst, weiß ich sicher, dass jeder Thread bis zu 4 Lösungen finden kann, aber wie kann ich dann zurück zum Host kopieren?

Ich übergebe ein riesiges Array mit genug Platz, um alle Threads speichern 4 Lösungen (4 Doppel für jede Lösung), und ein anderes mit der Anzahl der Lösungen pro Thread, aber das ist eine naive Lösung und ist der aktuelle Flaschenhals meines Kernels.

Ich mag es wirklich, dies zu optimieren. Das Hauptproblem besteht darin, eine variable Anzahl von Lösungen pro Thread in einem einzelnen Array zu verketten.

Answer 1

Die von Ihnen gesuchte Funktionalität wird als Stream-Komprimierung bezeichnet.

Sie müssen wahrscheinlich ein Array bereitstellen, das Platz für 4 Lösungen pro Thread enthält, da der Versuch, die Ergebnisse in einer kompakten Form zu speichern, wahrscheinlich so viele Abhängigkeiten zwischen den Threads erzeugt, die die Leistung beim Kopieren erhalten hat Weniger Daten zurück zum Host gehen durch eine längere Kernel-Ausführungszeit verloren. Die Ausnahme ist, wenn fast alle Threads keine Lösungen finden. In diesem Fall können Sie möglicherweise eine atomare Operation verwenden, um einen Index in einem Array zu verwalten. Für jede gefundene Lösung würden Sie sie also in einem Array an einem Index speichern und dann mit einer atomaren Operation den Index erhöhen. Ich denke, es wäre sicher atomicAdd() dafür zu verwenden. Vor dem Speichern eines Ergebnisses würde der Thread atomicAdd() verwenden, um den Index um eins zu erhöhen. atomicAdd() gibt den alten Wert zurück und der Thread kann das Ergebnis mit dem alten Wert als Index speichern.

In einer häufigeren Situation, in der es jedoch viele Ergebnisse gibt, ist die beste Lösung, einen Verdichtungsvorgang als separaten Schritt durchzuführen. Ein Weg, dies zu tun ist mit thrust::copy_if. Weitere Hintergrundinformationen finden Sie unter this question.