Conflict_Free_Offset-Makro, das im parallelen Präfixalgorithmus von GPU-Edelsteinen 3 verwendet wird

Question

Zunächst einmal ist der Link zum Algorithmus:

GPU-GEMS 3, Kapitel 39: Parallelpräfixsumme (Scan) mit CUDA .

Um Bankenkonflikte zu vermeiden, wird die Polsterung dem gemeinsam genutzten Speicherangebot jeder NUM_BANKS (d. H. 32 für Geräte von Ccccessing 2.x) Elemente hinzugefügt.Dies geschieht durch (wie in Abbildung 39-5):

generasacodicetagpre.

Ich verstehe nicht, wie AI / Num_BANKS dem Makro entspricht:

generasacodicetagpre.

ist es nicht gleich

generasacodicetagpre.

Jede Hilfe wird geschätzt.Danke

Answer 1

Ich habe diesen Code geschrieben und den Artikel mitbeschrieben, und ich bitte, dass Sie den Artikel nur zum Erlernen von Scanalgorithmen verwenden und den Code nicht verwenden. Es wurde geschrieben, als Cuda neu war, und ich war neu in Cuda. Wenn Sie eine moderne Implementierung des Scans in CUDA verwenden, benötigen Sie keine Bankkonfliktvermeidung.

Wenn Sie Scans den einfachen Weg tun möchten, verwenden Sie thrust::inclusive_scan oder thrust::exclusive_scan .

Wenn Sie einen Scan wirklich implementieren möchten, beziehen Sie sich auf aktuellere Artikel, z. B. diesen ein [1] . Oder für einen echten OPUS mit schnellerem Code, aber das erfordert ein bisschen mehr Studium, dieses ist [2] . Oder lesen Sie Sean Baxters Tutorial (obwohl letztere nicht Zitate der Samenarbeit auf den Scan-Algorithmus beinhaltet).

[1] Shubhabrata Sengupta, Mark Harris, Michael Girlande und John D. Owens. "Effiziente Parallel-Scan-Algorithmen für Vielkern-GPUs". In Jakub Kurzak, David A. BADER, und Jack Dongarra, Redakteure, wissenschaftliches Computer mit Multicore und Beschleunigten, Chapman & Hall / CRC Computational Science, Kapitel 19, Seiten 413-442. Taylor & Francis, Januar 2011. http://www.idav.ucdavis.edu/publications/print_pub?pub_id= 1041

[2] Merrill, D. und Grimshaw, A. Paralleler Scan für Streamarchitekturen. Technischer Bericht CS2009-14, Informatikabteilung, Universität von Virginia. Dezember 2009.