Konstante Speicher Auslastung in CUDA-code
Ich kann es nicht herausfinden mich, was ist der beste Weg um sicherzustellen, dass der Speicher in meinem kernel ist konstant. Es ist eine ähnliche Frage wie bei http://stackoverflow...r-pleasant-way.
Ich arbeite mit GTX 580 und kompilieren, nur für 2.0-Fähigkeit. Mein Kern sieht aus wie
__global__ Foo(const int *src, float *result) {...}Führe ich folgenden code auf dem host:
cudaMalloc(src, size);
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice);
Foo<<<...>>>(src, result);die alternative Möglichkeit ist das hinzufügen von
__constant__ src[size];zu .cu-Datei, entfernen Sie src Zeiger aus dem kernel und ausführen
cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice);
Foo<<<...>>>(result);Sind diese beiden Arten gleichwertig oder der erste übernimmt keine Garantie für die Verwendung der Konstanten-Speicher anstelle von globalen Speicher? Größe ändert sich dynamisch, so dass die zweite Methode ist nicht praktisch in meinem Fall.
Schreibe einen Kommentar
Du musst angemeldet sein, um einen Kommentar abzugeben.
Der zweite Weg ist der einzige Weg, um sicherzustellen, dass das array erstellt, um CUDA-Konstanten-Speicher und der Zugriff korrekt über die constant memory-cache. Aber sollten Sie sich Fragen, wie die Inhalte des Arrays gehen, auf die zugegriffen werden in einem block von threads. Wenn jeder thread Zugriff auf das array gleichmäßig, dann gibt es einen performance-Vorteil in der Verwendung von Konstanten-Speicher, weil es ist ein broadcast-Mechanismus, der von der Konstanten-Speicher-cache (das spart auch globalen Speicher-Bandbreite, weil eine ständige Erinnerung gespeichert ist, in offchip-DRAM-und cache reduziert die DRAM-Transaktion zählen). Aber wenn der Zugriff zufällig ist, dann kann es sein, serialisieren der Zugriff auf den lokalen Speicher, die negativ auf die Leistung.
Typische Dinge, die vielleicht gut passt für
__constant__Speicher wäre Koeffizienten Modell, GEWICHTE und andere Konstante Werte, die eingestellt werden müssen, um zur Laufzeit. Auf Fermi-GPUs, kernel-argument-Liste gespeichert ist, in steter Erinnerung, zum Beispiel. Wenn aber die Inhalte sind der Zugang nicht einheitlich, und der Art oder Größe der Mitglieder ist nicht konstant aufrufen, dann ganz normal globalen Speicher vorzuziehen ist.Beachten Sie auch, dass es ein limit von 64 Kb Konstanten-Speicher pro GPU-Kontext, so ist es nicht praktisch zum speichern von sehr großen Datenmengen in steter Erinnerung. Wenn Sie eine Menge von nur-lese-Speicher mit cache, es könnte einen Versuch Wert sein die Bindung der Daten an eine textur und sehen, wie die performance ist. Auf pre-Fermi-Karten, es liefert meist eine praktische Leistung zu gewinnen, die auf der Fermi -, können die Ergebnisse weniger vorhersehbar im Vergleich zum globalen Speicher wegen der Verbesserung der cache-layout, dass in Architektur.
Die Erste Methode garantiert die Speicher konstant ist innerhalb der Funktion
Foo. Die beiden sind nicht gleichwertig, die zweite garantiert es Konstante es nach der Initialisierung. Wenn Sie brauchen dynamische, als Sie nee, etwas zu verwenden, ähnlich wie der erste Weg.